sync 2e7068faf55e

Browse files

Files changed (7) hide show

README.md +76 -0
build/webgpu/bench.json +114 -0
build/webgpu/manifest.json +188 -0
build/webgpu/metadata.json +19 -0
build/webgpu/pad-quad-vec4.wgsl.jinja +113 -0
build/webgpu/pad.wgsl.jinja +100 -0
build/webgpu/test.json +1383 -0

README.md CHANGED Viewed

@@ -1,3 +1,79 @@
 ---
 license: apache-2.0
 ---

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.Pad
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 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.
+See the [ONNX `Pad` spec](https://onnx.ai/onnx/operators/onnx__Pad.html) for the reference semantics.
+## 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
+Attributes and default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `mode` | `"constant"` | Padding mode: `constant` (default), `reflect`, `edge`, or `wrap`. |
+| `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]`. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16`, `uint32`, `int32`, `uint8`, `int8`, `bool` |
+## 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
+- [`pad-quad-vec4.wgsl.jinja`](build/webgpu/pad-quad-vec4.wgsl.jinja)
+- [`pad.wgsl.jinja`](build/webgpu/pad.wgsl.jinja)
+## 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.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/ai.onnx.Pad", { version: 1 });
+// Explicit destinations request optional results or supply metadata that cannot be inferred.
+const { output } = await kernel({ data: { data: dataData, shape: [2] } }, {
+  attrs: { pads: [0, 0] },
+  outputs: { output: { shape: [2], dtype: "float32" } },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,114 @@

+{
+  "op": "ai.onnx.Pad",
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "rank4_image",
+      "attrs": { "pads": [0, 0, 1, 1, 0, 0, 1, 1] },
+      "inputs": { "data": { "dtype": "float32", "shape": [1, 16, 128, 128], "dist": "normal", "seed": 1 } },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 16, 130, 130], "dist": "empty" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.data) + numel(shapes.output)) * 4" }]
+      }
+    },
+    {
+      "name": "rank4_image_inner_aligned",
+      "attrs": { "pads": [0, 0, 1, 1, 0, 0, 1, 1] },
+      "inputs": { "data": { "dtype": "float32", "shape": [1, 16, 126, 126], "dist": "normal", "seed": 2 } },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 16, 128, 128], "dist": "empty" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.data) + numel(shapes.output)) * 4" }]
+      },
+      "preset": "smoke"
+    },
+    {
+      "name": "rank3_seq_pad_llm",
+      "attrs": { "pads": [0, 0, 0, 0, 64, 0] },
+      "inputs": { "data": { "dtype": "float32", "shape": [1, 448, 768], "dist": "normal", "seed": 3 } },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 512, 768], "dist": "empty" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.data) + numel(shapes.output)) * 4" }]
+      }
+    },
+    {
+      "name": "rank4_image_reflect_inner_aligned",
+      "attrs": { "pads": [0, 0, 2, 2, 0, 0, 2, 2], "mode": "reflect" },
+      "inputs": { "data": { "dtype": "float32", "shape": [1, 16, 128, 128], "dist": "normal", "seed": 4 } },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 16, 132, 132], "dist": "empty" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.data) + numel(shapes.output)) * 4" }]
+      }
+    },
+    {
+      "name": "rank4_image_f16_inner_aligned",
+      "attrs": { "pads": [0, 0, 1, 1, 0, 0, 1, 1] },
+      "inputs": { "data": { "dtype": "float16", "shape": [1, 16, 126, 126], "dist": "normal", "seed": 5 } },
+      "outputs": { "output": { "dtype": "float16", "shape": [1, 16, 128, 128], "dist": "empty" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.data) + numel(shapes.output)) * 2" }]
+      }
+    },
+    {
+      "name": "rank4_image_value_input_inner_aligned",
+      "attrs": { "pads": [0, 0, 1, 1, 0, 0, 1, 1] },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [1, 16, 126, 126], "dist": "normal", "seed": 6 },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 16, 128, 128], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.data) + numel(shapes.output)) * 4" }] }
+    },
+    {
+      "name": "rank4_image_edge_inner_aligned",
+      "attrs": { "pads": [0, 0, 1, 1, 0, 0, 1, 1], "mode": "edge" },
+      "inputs": { "data": { "dtype": "float32", "shape": [1, 16, 126, 126], "dist": "normal", "seed": 7 } },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 16, 128, 128], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.data) + numel(shapes.output)) * 4" }] }
+    },
+    {
+      "name": "rank4_image_f16_reflect_inner_aligned",
+      "attrs": { "pads": [0, 0, 2, 2, 0, 0, 2, 2], "mode": "reflect" },
+      "inputs": { "data": { "dtype": "float16", "shape": [1, 16, 128, 128], "dist": "normal", "seed": 8 } },
+      "outputs": { "output": { "dtype": "float16", "shape": [1, 16, 132, 132], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.data) + numel(shapes.output)) * 2" }] }
+    },
+    {
+      "name": "rank3_value_input_forces_scalar_path_large",
+      "preset": "stress",
+      "attrs": { "mode": "constant", "pads": [0, 0, 0, 0, 64, 0] },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [8, 448, 768], "dist": "normal", "seed": 11, "scale": 2 },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [8, 512, 768], "dist": "empty" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.data) + numel(shapes.output)) * 4" }]
+      }
+    },
+    {
+      "name": "rank4_edge_inner_unaligned_forces_scalar_path",
+      "preset": "stress",
+      "attrs": { "mode": "edge", "pads": [0, 0, 1, 1, 0, 0, 1, 1] },
+      "inputs": { "data": { "dtype": "float32", "shape": [1, 64, 127, 127], "dist": "normal", "seed": 12, "scale": 2 } },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 64, 129, 129], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.data) + numel(shapes.output)) * 4" }] }
+    },
+    {
+      "name": "rank2_constant_inner_unaligned_forces_scalar_path",
+      "preset": "stress",
+      "attrs": { "mode": "constant", "pads": [1, 1, 1, 0] },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [2046, 1024], "dist": "normal", "seed": 13, "scale": 2 },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [3.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2048, 1025], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "(numel(shapes.data) + numel(shapes.output)) * 4" }] }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,188 @@

+{
+  "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"
+            }
+          },
+          "bindings": "valueInput",
+          "dispatch": { "threads": "numel(shapes.output)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "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
+            }
+          },
+          "bindings": "scalarPad",
+          "dispatch": { "threads": "numel(shapes.output)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,19 @@

+{
+  "name": "ai.onnx.Pad",
+  "id": "_ai_onnx_pad_webgpu_f677a1e",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "c5OFqXWDpHEMBxvQhbwziuREKzcvtHq7AD3mAWc/O/w=",
+      "manifest.json": "LdYMiDSOmaf3p/AM+35OYK9r6K9ujcB5KaVJwyW/gWc=",
+      "pad-quad-vec4.wgsl.jinja": "1CGMCSmy+Suuii9slZko0MXqHiIf0+iAb8UjG/mKLkU=",
+      "pad.wgsl.jinja": "/dyKjV18AkgaduRKI8SGVZz6qhpg3i/XVAj4IROc5r8=",
+      "test.json": "9NikPawZjt0OZY+069sKOCAJ/K5Xc7F6usOSJ6TrvpE="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Pad" }
+}

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

	@@ -0,0 +1,113 @@

+// Vec4 form for outputs whose innermost dimension is four-aligned. One invocation
+// handles four consecutive innermost elements and decomposes the outer
+// coordinates once; only the innermost axis runs map_coord per element. This
+// supports constant, edge, reflect, and wrap modes.
+//
+// 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;
+  }
+{% if mode == "constant" %}
+  return -1i;
+{% elif mode == "edge" %}
+  if (dim <= 0i) {
+    return -1i;
+  }
+  return min(max(coord, 0i), dim - 1i);
+{% elif mode == "wrap" %}
+  if (dim <= 0i) {
+    return -1i;
+  }
+  return ((coord % dim) + dim) % dim;
+{% else %}
+  if (dim < 2i) {
+    return -1i;
+  }
+  let period = 2i * dim - 2i;
+  let wrapped = ((coord % period) + period) % period;
+  if (wrapped < dim) {
+    return wrapped;
+  }
+  return period - wrapped;
+{% endif %}
+}
+{% 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;
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+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);
+{% if rank > 1 %}
+    var rem = v / OUT_W_VECS;
+{% endif %}
+    var srcBase = 0i;
+    var outerValid = true;
+{% 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);
+    rem = rem % {{ safeOStride }}u;
+    let mapped{{ axis }} = map_coord(out_coord{{ axis }} - {{ padBefore }}i, {{ sliced }}i);
+    outerValid = outerValid && (mapped{{ axis }} >= 0i);
+    srcBase += (mapped{{ axis }} + {{ crop }}i) * {{ dstride.value }}i;
+{% endfor %}
+    var result = {{ vectorScalar }}(fill);
+    if (outerValid) {
+      let m0 = map_coord(ow0 - PAD_W, SLICED_W);
+      let m1 = map_coord(ow0 + 1i - PAD_W, SLICED_W);
+      let m2 = map_coord(ow0 + 2i - PAD_W, SLICED_W);
+      let m3 = map_coord(ow0 + 3i - PAD_W, SLICED_W);
+      if (m0 >= 0i) { result.x = data[u32(srcBase + m0 + CROP_W)]; }
+      if (m1 >= 0i) { result.y = data[u32(srcBase + m1 + CROP_W)]; }
+      if (m2 >= 0i) { result.z = data[u32(srcBase + m2 + CROP_W)]; }
+      if (m3 >= 0i) { result.w = data[u32(srcBase + m3 + CROP_W)]; }
+    }
+    output[v] = result;
+  }
+}

build/webgpu/pad.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,100 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+fn map_coord(coord: i32, dim: i32) -> i32 {
+  if (coord >= 0i && coord < dim) {
+    return coord;
+  }
+{% if mode == "constant" %}
+  return -1i;
+{% elif mode == "edge" %}
+  if (dim <= 0i) {
+    return -1i;
+  }
+  return min(max(coord, 0i), dim - 1i);
+{% elif mode == "wrap" %}
+  if (dim <= 0i) {
+    return -1i;
+  }
+  return ((coord % dim) + dim) % dim;
+{% else %}
+  if (dim < 2i) {
+    return -1i;
+  }
+  let period = 2i * dim - 2i;
+  let wrapped = ((coord % period) + period) % period;
+  if (wrapped < dim) {
+    return wrapped;
+  }
+  return period - wrapped;
+{% endif %}
+}
+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 %}
+{% set crop_start = 0 %}
+{% endif %}
+{% if post < 0 %}
+{% set crop_end_trim = 0 - post %}
+{% 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 %}
+{% endif %}
+{% if pre > 0 %}
+{% set pad_before = pre %}
+{% else %}
+{% set pad_before = 0 %}
+{% endif %}
+  let out_coord{{ axis }} = i32(rem / {{ out_stride_divisor }}u);
+  rem = rem % {{ out_stride_divisor }}u;
+  let mapped_coord{{ axis }} = map_coord(out_coord{{ axis }} - {{ pad_before }}i, {{ sliced_dim }}i);
+  if (mapped_coord{{ axis }} < 0i) {
+    return -1i;
+  }
+  let in_coord{{ axis }} = mapped_coord{{ axis }} + {{ crop_start }}i;
+  offset = offset + in_coord{{ axis }} * {{ data_stride.value }}i;
+{% endfor %}
+  return offset;
+}
+@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 past the
+  // maxComputeWorkgroupsPerDimension limit.
+  let i = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if (i >= params.count) {
+    return;
+  }
+  let src = data_offset_or_invalid(i);
+  if (src < 0i) {
+{% if valueFromInput %}
+    output[i] = padValue[0];
+{% else %}
+    output[i] = {{ scalar }}({{ source.value }});
+{% endif %}
+  } else {
+    output[i] = data[u32(src)];
+  }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,1383 @@

+{
+  "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]
+  },
+  "cases": [
+    {
+      "name": "rank2_constant",
+      "attrs": { "pads": [1, 2, 1, 0] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 5] } }
+    },
+    {
+      "name": "rank4_f16",
+      "attrs": { "pads": [0, 0, 1, 1, 0, 0, 1, 1] },
+      "inputs": { "data": { "dtype": "float16", "shape": [1, 1, 2, 2] } },
+      "outputs": { "output": { "dtype": "float16", "shape": [1, 1, 4, 4] } },
+      "tolerance": 0.001
+    },
+    {
+      "name": "ort_value_input_constant_1d",
+      "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": {
+        "data": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.0, 2.0] } },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [123.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [5], "tolerance": 0 } }
+    },
+    {
+      "name": "value_input_float16_constant_1d",
+      "attrs": { "mode": "constant", "pads": [1, 2] },
+      "inputs": {
+        "data": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [1.5, -2.0] } },
+        "constant_value": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [-3.5] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float16",
+          "shape": [5],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [-3.5, 1.5, -2.0, -3.5, -3.5] }
+        }
+      }
+    },
+    {
+      "name": "value_input_int8_constant_1d",
+      "attrs": { "mode": "constant", "pads": [1, 1] },
+      "inputs": {
+        "data": { "dtype": "int8", "shape": [2], "data": { "kind": "values", "values": [-128, 127] } },
+        "constant_value": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [-7] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "int8",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [-7, -128, 127, -7] }
+        }
+      }
+    },
+    {
+      "name": "value_input_uint8_constant_1d",
+      "attrs": { "mode": "constant", "pads": [1, 1] },
+      "inputs": {
+        "data": { "dtype": "uint8", "shape": [2], "data": { "kind": "values", "values": [0, 255] } },
+        "constant_value": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [173] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "uint8",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [173, 0, 255, 173] }
+        }
+      }
+    },
+    {
+      "name": "value_input_bool_constant_1d",
+      "attrs": { "mode": "constant", "pads": [1, 1] },
+      "inputs": {
+        "data": { "dtype": "bool", "shape": [2], "data": { "kind": "values", "values": [0, 1] } },
+        "constant_value": { "dtype": "bool", "shape": [], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "bool",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [1, 0, 1, 1] }
+        }
+      }
+    },
+    {
+      "name": "ort_int32_value_input_constant_1d",
+      "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": {
+        "data": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [1, 2] } },
+        "constant_value": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [123] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [5], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_bool_constant_pad_axis1",
+      "provenance": { "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc", "test": "PadOpTest.BoolType" },
+      "attrs": { "mode": "constant", "pads": [0, 2, 0, 0] },
+      "inputs": {
+        "data": { "dtype": "bool", "shape": [3, 2], "data": { "kind": "values", "values": [1, 0, 1, 0, 1, 0] } },
+        "constant_value": { "dtype": "bool", "shape": [], "data": { "kind": "values", "values": [1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "bool",
+          "shape": [3, 4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [1, 1, 1, 0, 1, 1, 1, 0, 1, 1, 1, 0] }
+        }
+      }
+    },
+    {
+      "name": "ort_constant_pad_axes_out_of_order_rank4",
+      "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": {
+        "data": {
+          "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": {
+        "output": {
+          "dtype": "float32",
+          "shape": [1, 2, 2, 4],
+          "tolerance": 0,
+          "data": {
+            "kind": "values",
+            "values": [0.0, 1.0, 1.0, 0.0, 0.0, 1.0, 1.0, 0.0, 0.0, 1.0, 1.0, 0.0, 0.0, 1.0, 1.0, 0.0]
+          }
+        }
+      }
+    },
+    {
+      "name": "ort_constant_pad_axes_int32_axis1_axis3",
+      "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": {
+        "data": {
+          "dtype": "int32",
+          "shape": [1, 2, 2, 2],
+          "data": { "kind": "values", "values": [1, 1, 1, 1, 1, 1, 1, 1] }
+        }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [1, 2, 2, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_constant_pad_axes_last_two_both_sides",
+      "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": {
+        "data": {
+          "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": { "output": { "dtype": "float32", "shape": [1, 2, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_constant_pad_full_rank_without_axes",
+      "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": {
+        "data": {
+          "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": { "output": { "dtype": "float32", "shape": [1, 2, 2, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "rank5_constant",
+      "attrs": { "pads": [0, 1, 0, 1, 0, 0, 0, 1, 0, 1] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 2, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.19 }
+        },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-2.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 3, 2, 3, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank6_constant",
+      "attrs": { "pads": [0, 1, 0, 0, 1, 0, 0, 0, 1, 0, 0, 1] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 2, 1, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.19 }
+        },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [7.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 3, 2, 2, 2, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "backend_spec_example_axis1_zero",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_Spec_Example"
+      },
+      "attrs": { "pads": [0, 2, 0, 0] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [3, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_constant_1d",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_Constant_1D"
+      },
+      "attrs": { "pads": [1, 2] },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.0, 2.0] } },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [123.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [5], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_constant_1d_zero_pads",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_Constant_1D_Zero"
+      },
+      "attrs": { "pads": [0, 0] },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.0, 2.0] } },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [123.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_constant_2d",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_Constant_2D"
+      },
+      "attrs": { "pads": [1, 2, 1, 2] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 2],
+          "data": { "kind": "values", "values": [11.0, 21.0, 12.0, 22.0] }
+        },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [123.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 6], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_edge_2d_axis1",
+      "provenance": { "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc", "test": "PadOpTest.Pad_Edge_1D" },
+      "attrs": { "pads": [0, 2, 0, 1], "mode": "edge" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [3, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 5], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_reflect_2d_axis1",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_Reflect_1D"
+      },
+      "attrs": { "pads": [0, 1, 0, 1], "mode": "reflect" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [3, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_wrap_2d_axis1",
+      "provenance": { "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc", "test": "PadOpTest.Pad_Wrap_1D" },
+      "attrs": { "pads": [0, 1, 0, 1], "mode": "wrap" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [3, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_wrap_1d_upper_pad_greater_than_input",
+      "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": {
+        "data": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [8],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 1.0, 2.0, 3.0, 1.0, 2.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_wrap_1d_lower_pad_greater_than_input",
+      "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": {
+        "data": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [8],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [2.0, 3.0, 1.0, 2.0, 3.0, 1.0, 2.0, 3.0] }
+        }
+      }
+    },
+    {
+      "name": "ort_edge_2d_full_pad",
+      "provenance": { "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc", "test": "PadOpTest.Pad_Edge_2D" },
+      "attrs": { "pads": [2, 2, 2, 2], "mode": "edge" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [11.0, 21.0, 31.0, 12.0, 22.0, 32.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [6, 7], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_reflect_2d_full_pad",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_Reflect_2D"
+      },
+      "attrs": { "pads": [2, 2, 2, 2], "mode": "reflect" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [11.0, 21.0, 31.0, 12.0, 22.0, 32.0, 13.0, 23.0, 33.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [7, 7], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_wrap_2d_full_pad",
+      "provenance": { "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc", "test": "PadOpTest.Pad_Wrap_2D" },
+      "attrs": { "pads": [2, 2, 2, 2], "mode": "wrap" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [11.0, 21.0, 31.0, 12.0, 22.0, 32.0, 13.0, 23.0, 33.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [7, 7], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_wrap_negative_front_positive_back",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_Wrap_NegativeFront_PositiveBack"
+      },
+      "attrs": { "pads": [-3, 3], "mode": "wrap" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_reflect_pad_equals_extent_minus_one",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_Reflect_PadEqualsExtentMinus1_Succeeds"
+      },
+      "attrs": { "pads": [2, 2], "mode": "reflect" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [7], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "constant_1d_negative_front_crop_post_pad",
+      "attrs": { "pads": [-1, 2] },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0] } },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [123.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [5], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "constant_2d_negative_axis0_crop",
+      "attrs": { "pads": [-1, 0, -1, 0] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [3, 3],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0] }
+        },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-5.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "constant_vec4_negative_trailing_crop_both_axes",
+      "provenance": {
+        "source": "ONNX Pad-19 specification",
+        "test": "negative end pads crop both axes on the packed default-value path"
+      },
+      "attrs": { "pads": [0, 0, -1, -1] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [3, 5],
+          "data": {
+            "kind": "values",
+            "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0]
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 4], "tolerance": 0 } }
+    },
+    {
+      "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 } }
+    },
+    {
+      "name": "ort_constant_2d_negative_back_pad_crop",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_Constant_2D_negative_pads_1"
+      },
+      "attrs": { "pads": [1, 2, 1, -1] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [11.0, 21.0, 31.0, 12.0, 22.0, 32.0] }
+        },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [123.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_constant_3d_complex",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_3D_complex"
+      },
+      "attrs": { "pads": [1, 0, 0, -1, 0, 0] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 2, 2],
+          "data": { "kind": "values", "values": [11.0, 12.0, 21.0, 22.0, 111.0, 112.0, 121.0, 122.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "constant_3d_mixed_signs_crop_and_pad",
+      "attrs": { "pads": [0, -1, 1, 0, 1, -2] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 3, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 1.0, 2.0, 3.0, 10.0, 11.0, 12.0, 13.0, 20.0, 21.0, 22.0, 23.0, 100.0, 101.0, 102.0, 103.0, 110.0, 111.0, 112.0, 113.0, 120.0, 121.0, 122.0, 123.0]
+          }
+        },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [9.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_constant_4d_crop_removes_all_source_values",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.ConstantFill_F32_RemovesAllDataOnAxis"
+      },
+      "attrs": { "pads": [0, 0, -4, 0, 0, 0, 4, 0] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [1, 1, 4, 4],
+          "data": {
+            "kind": "values",
+            "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0]
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 1, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_constant_3d_innermost_crop_then_post_pad",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.ConstantMode_InnermostCropThenPostPad"
+      },
+      "attrs": { "pads": [1, 3, -2, -1, 0, 1] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 3, 5],
+          "data": {
+            "kind": "values",
+            "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0, 9.0, 10.0, 11.0, 12.0, 13.0, 14.0, 15.0, 16.0, 17.0, 18.0, 19.0, 20.0, 21.0, 22.0, 23.0, 24.0, 25.0, 26.0, 27.0, 28.0, 29.0, 30.0]
+          }
+        },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [9.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 6, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_constant_3d_mixed_signs_small",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.ConstantMode_MixedSigns_Small"
+      },
+      "attrs": { "pads": [1, 3, -2, -1, 0, 1] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 6, 4],
+          "data": { "kind": "cycle", "values": [1.0, 2.0, 3.0, 4.0, 5.0] }
+        },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [9.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 9, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_edge_mode_extent_one_after_crop",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.EdgeMode_ExtentOne_Valid"
+      },
+      "attrs": { "pads": [-3, 3], "mode": "edge" },
+      "inputs": { "data": { "dtype": "float32", "shape": [4], "data": { "kind": "constant", "value": 1.0 } } },
+      "outputs": { "output": { "dtype": "float32", "shape": [4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_edge_mode_flattened_innermost_axis",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.EdgeMode_FlattenedInnermostAxis"
+      },
+      "attrs": { "pads": [0, 0, 0, 0, 0, 0, 0, 1], "mode": "edge" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [2, 3, 2, 4], "data": { "kind": "linspace", "start": 0.0, "end": 47.0 } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 2, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_zero_dim_1d_pad_to_nonempty",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_Constant_DimWithZeroInput / 1D"
+      },
+      "attrs": { "pads": [1, 1] },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_zero_dim_1d_empty_pads",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_Constant_DimWithZeroInput / 1D empty pads"
+      },
+      "attrs": { "pads": [0, 0] },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [0], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_zero_dim_2d_inner_pad_to_nonempty",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_Constant_DimWithZeroInput / 2D"
+      },
+      "attrs": { "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [2, 0], "data": { "kind": "values", "values": [] } },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_zero_dim_3d_middle_pad_to_nonempty_f16",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_Constant_DimWithZeroInput / 3D"
+      },
+      "attrs": { "pads": [0, 1, 0, 0, 1, 0] },
+      "inputs": {
+        "data": { "dtype": "float16", "shape": [2, 0, 2], "data": { "kind": "values", "values": [] } },
+        "constant_value": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [2, 2, 2], "tolerance": 0.001 } }
+    },
+    {
+      "name": "ort_edge_zero_dim_unpadded_axis_stays_empty",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_Edge_DimWithZeroInput",
+        "notes": "Valid ORT edge-mode case where the zero-sized axis is not padded, so the output remains empty."
+      },
+      "attrs": { "mode": "edge", "pads": [1, 0, 1, 0] },
+      "inputs": { "data": { "dtype": "float32", "shape": [2, 0], "data": { "kind": "values", "values": [] } } },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 0], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_reflect_zero_dim_unpadded_axis_stays_empty",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_Reflect_DimWithZeroInput",
+        "notes": "Valid ORT reflect-mode case where the zero-sized axis is not padded, so the output remains empty."
+      },
+      "attrs": { "mode": "reflect", "pads": [1, 0, 1, 0] },
+      "inputs": { "data": { "dtype": "float32", "shape": [2, 0], "data": { "kind": "values", "values": [] } } },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 0], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_constant_2d_crop_front_axis0",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_Constant_2D_negative_pads_2"
+      },
+      "attrs": { "pads": [-1, 0, 0, 0], "mode": "constant" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [11.0, 21.0, 31.0, 12.0, 22.0, 32.0] }
+        },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [123.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_constant_3d_crop_both_last_axis",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_Constant_3D_negative_pads"
+      },
+      "attrs": { "pads": [0, 0, -1, 0, 0, -1], "mode": "constant" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [1, 1, 3], "data": { "kind": "values", "values": [0.0, 1.0, 2.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 1, 1], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_constant_4d_negative_pads_inner_window",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_Constant_4D_negative_pads"
+      },
+      "attrs": { "pads": [0, 0, -1, -3, 0, 0, -2, -4], "mode": "constant" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [1, 1, 10, 10],
+          "data": { "kind": "linspace", "start": 0.0, "end": 99.0 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 1, 7, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_constant_fill_removes_all_data_on_axis",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.ConstantFill_F32_RemovesAllDataOnAxis"
+      },
+      "attrs": { "pads": [0, 0, -4, 0, 0, 0, 4, 0], "mode": "constant" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [1, 1, 4, 4], "data": { "kind": "linspace", "start": 1.0, "end": 16.0 } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 1, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_constant_large_negative_pad_no_output",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.ConstantPadLargeNegativePadNoOutput",
+        "notes": "Large positive secondary dimension is paired with a zero leading dimension, so the output has zero elements."
+      },
+      "attrs": { "pads": [1, 1048576, -2, -3, 0, 1], "mode": "constant" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 18, 4],
+          "data": { "kind": "cycle", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [100.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [0, 1048594, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_constant_mixed_signs_small",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.ConstantMode_MixedSigns_Small"
+      },
+      "attrs": { "pads": [1, 3, -2, -1, 0, 1], "mode": "constant" },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 6, 4],
+          "data": { "kind": "cycle", "values": [1.0, 2.0, 3.0, 4.0, 5.0] }
+        },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [9.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 9, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_constant_innermost_crop_then_post_pad",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.ConstantMode_InnermostCropThenPostPad"
+      },
+      "attrs": { "pads": [1, 3, -2, -1, 0, 1], "mode": "constant" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 1.0, "end": 30.0 } },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [9.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 6, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_edge_extent_one_after_crop",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.EdgeMode_ExtentOne_Valid"
+      },
+      "attrs": { "pads": [-3, 3], "mode": "edge" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.0, 1.0, 1.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_constant_3d_complex_crop_and_prepad",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_3D_complex"
+      },
+      "attrs": { "pads": [1, 0, 0, -1, 0, 0] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 2, 2],
+          "data": { "kind": "values", "values": [11.0, 12.0, 21.0, 22.0, 111.0, 112.0, 121.0, 122.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_reflect_3d_inner_no_padding",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_Reflect_3D_Inner_No_Padding"
+      },
+      "attrs": { "pads": [1, 1, 0, 1, 1, 0], "mode": "reflect" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [3, 2, 5], "data": { "kind": "linspace", "start": 1.0, "end": 30.0 } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [5, 4, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_wrap_3d_inner_no_padding",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_wrap_3D_Inner_No_Padding"
+      },
+      "attrs": { "pads": [1, 1, 0, 1, 1, 0], "mode": "wrap" },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [3, 2, 5], "data": { "kind": "linspace", "start": 1.0, "end": 30.0 } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [5, 4, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_constant_pad",
+      "attrs": { "mode": "constant", "pads": [0, 0, 1, 3, 0, 0, 2, 4] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [1, 3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_constant_pad_input_data" } }
+        },
+        "constant_value": {
+          "dtype": "float32",
+          "shape": [],
+          "data": { "kind": "values", "values": [1.2000000476837158] }
+        }
+      },
+      "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."
+      }
+    },
+    {
+      "name": "onnx_backend_constant_pad_axes",
+      "attrs": { "mode": "constant", "pads": [0, 0, 0, 3, 0, 0, 0, 4] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [1, 3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_constant_pad_input_data" } }
+        },
+        "constant_value": {
+          "dtype": "float32",
+          "shape": [],
+          "data": { "kind": "values", "values": [1.2000000476837158] }
+        }
+      },
+      "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."
+      }
+    },
+    {
+      "name": "onnx_backend_edge_pad",
+      "attrs": { "mode": "edge", "pads": [0, 0, 1, 1, 0, 0, 1, 1] },
+      "inputs": {
+        "data": {
+          "dtype": "int32",
+          "shape": [1, 3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [1, 0, 0, 2, 1, 0, 0, 0, 0, 0, 0, 1, 0, 0, 0, 0, 1, 0, 0, 0, -2, 0, 0, 0, 2, -1, 0, 0, 1, 1, 0, 0, 0, -1, 0, 0, 1, 1, 0, 0, -1, -1, -1, 1, 0, 0, -1, 0, -1, 0, 0, 0, 0, -1, 0, 0, 0, 0, 0, 0]
+          }
+        }
+      },
+      "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."
+      }
+    },
+    {
+      "name": "onnx_backend_reflect_pad",
+      "attrs": { "mode": "reflect", "pads": [0, 0, 1, 1, 0, 0, 1, 1] },
+      "inputs": {
+        "data": {
+          "dtype": "int32",
+          "shape": [1, 3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [0, 0, 0, -1, 0, 0, -1, 0, 0, 0, 0, 0, 1, -1, 0, 0, 0, 0, 0, 0, -1, 0, 0, -1, 1, 1, 1, 0, -1, 1, 0, 1, 0, 0, 0, 0, 0, 1, 0, 0, 1, -1, -1, 0, -1, 1, 0, 0, 1, 1, 1, 0, 0, 1, 0, 0, 0, 0, 0, 0]
+          }
+        }
+      },
+      "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."
+      }
+    },
+    {
+      "name": "onnx_backend_wrap_pad",
+      "attrs": { "mode": "wrap", "pads": [0, 0, 1, 1, 0, 0, 1, 1] },
+      "inputs": {
+        "data": {
+          "dtype": "int32",
+          "shape": [1, 3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [0, -1, 0, 1, 0, 0, 0, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, -1, -1, 0, 0, 0, 2, 0, 0, 1, -1, 0, 0, 1, 0, 0, 0, 0, 1, -1, -1, 0, 0, 1, 0, 0, -1, 0, -1, -1, 1, 0, 0, 0, 0, 0, -1, 0, 0, 0, 0, 0]
+          }
+        }
+      },
+      "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."
+      }
+    },
+    {
+      "name": "ort_edge_3d",
+      "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": {
+        "data": {
+          "dtype": "float32",
+          "shape": [1, 2, 3],
+          "data": { "kind": "values", "values": [11.0, 21.0, 31.0, 12.0, 22.0, 32.0] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 6, 7], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_constant_3d_inner_no_padding",
+      "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": {
+        "data": { "dtype": "float32", "shape": [3, 2, 5], "data": { "kind": "linspace", "start": 1.0, "end": 30.0 } },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [31.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [5, 4, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_edge_3d_inner_no_padding",
+      "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": {
+        "data": { "dtype": "float32", "shape": [3, 2, 5], "data": { "kind": "linspace", "start": 1.0, "end": 30.0 } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [5, 4, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_edge_3d_last_pad_slice_inner_no_padding",
+      "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": {
+        "data": { "dtype": "float32", "shape": [3, 2, 5], "data": { "kind": "linspace", "start": 1.0, "end": 30.0 } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [5, 2, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_edge_3d_last_slice_inner_no_padding",
+      "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": {
+        "data": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 1.0, "end": 30.0 } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 2, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_reflect_3d_last_pad_slice_inner_no_padding",
+      "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": {
+        "data": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 1.0, "end": 30.0 } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 3, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_reflect_3d_last_slice_inner_no_padding",
+      "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": {
+        "data": { "dtype": "float32", "shape": [2, 3, 5], "data": { "kind": "linspace", "start": 1.0, "end": 30.0 } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 2, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_wrap_3d_inner_no_padding2",
+      "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": {
+        "data": { "dtype": "float32", "shape": [3, 2, 5], "data": { "kind": "linspace", "start": 1.0, "end": 30.0 } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [5, 6, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int8_constant_1d_edge_values",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_Constant_1D",
+        "notes": "Constant 1D padding with logical int8 edge values."
+      },
+      "attrs": { "mode": "constant", "pads": [2, 1] },
+      "inputs": {
+        "data": { "dtype": "int8", "shape": [3], "data": { "kind": "values", "values": [-128, 0, 127] } },
+        "constant_value": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [-1] } }
+      },
+      "outputs": { "output": { "dtype": "int8", "shape": [6], "tolerance": 0 } }
+    },
+    {
+      "name": "uint8_constant_crop_then_post_pad",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.ConstantMode_InnermostCropThenPostPad",
+        "notes": "Crops the first column, then post-pads with a uint8 edge value."
+      },
+      "attrs": { "mode": "constant", "pads": [0, -1, 0, 2] },
+      "inputs": {
+        "data": { "dtype": "uint8", "shape": [2, 3], "data": { "kind": "values", "values": [0, 1, 2, 253, 254, 255] } },
+        "constant_value": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [255] } }
+      },
+      "outputs": { "output": { "dtype": "uint8", "shape": [2, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "int32_constant_value_exact_below_negative_float24",
+      "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": {
+        "data": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [-123, 123] } },
+        "constant_value": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [-16777217] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "int32",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [-16777217, -123, 123, -16777217] }
+        }
+      }
+    },
+    {
+      "name": "constant_two_pass_rank4_border",
+      "attrs": { "mode": "constant", "pads": [0, 0, 1, 2, 0, 0, 1, 2] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [1, 2, 6, 8],
+          "data": { "kind": "fillFloat32", "sinStep": 0.17, "cosStep": 0.23 }
+        },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 8, 12], "tolerance": 0 } }
+    },
+    {
+      "name": "constant_two_pass_unaligned_interior_tail",
+      "attrs": { "mode": "constant", "pads": [1, 0, 1, 0, 1, 2] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 3, 5],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.29 }
+        },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [9.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 4, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "constant_two_pass_crop_and_pad",
+      "attrs": { "mode": "constant", "pads": [1, 0, -2, 0, 1, -4] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 3, 10],
+          "data": { "kind": "fillFloat32", "sinStep": 0.19, "cosStep": 0.13 }
+        },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [7.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "constant_two_pass_rank1_crop_unaligned",
+      "attrs": { "mode": "constant", "pads": [3, -1] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.31, "cosStep": 0.07 }
+        },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-2.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [8], "tolerance": 0 } }
+    },
+    {
+      "name": "constant_two_pass_f16_border",
+      "attrs": { "mode": "constant", "pads": [0, 0, 1, 1, 0, 0, 1, 1] },
+      "inputs": {
+        "data": { "dtype": "float16", "shape": [1, 1, 3, 6] },
+        "constant_value": { "dtype": "float16", "shape": [], "data": { "kind": "values", "values": [0.5] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [1, 1, 5, 8] } },
+      "tolerance": 0.001
+    },
+    {
+      "name": "edge_3d_quad_inner_aligned",
+      "attrs": { "mode": "edge", "pads": [1, 1, 1, 1, 1, 1] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 3, 6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.21, "cosStep": 0.17 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 5, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "reflect_3d_quad_inner_aligned",
+      "attrs": { "mode": "reflect", "pads": [1, 1, 3, 1, 1, 3] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 3, 6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.27, "cosStep": 0.19 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 5, 12], "tolerance": 0 } }
+    },
+    {
+      "name": "wrap_3d_quad_inner_aligned",
+      "attrs": { "mode": "wrap", "pads": [1, 1, 4, 1, 1, 2] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [2, 3, 6],
+          "data": { "kind": "fillFloat32", "sinStep": 0.33, "cosStep": 0.11 }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 5, 12], "tolerance": 0 } }
+    },
+    {
+      "name": "int32_constant_attr_value_exact_above_float24_scalar_path_unaligned",
+      "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": {
+        "data": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [7] } },
+        "constant_value": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [16777217] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "int32",
+          "shape": [3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [16777217, 7, 16777217] }
+        }
+      }
+    },
+    {
+      "name": "uint32_value_input_exact_above_float24",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_Constant_1D",
+        "notes": "An optional uint32 value input above 2^24 must remain exact rather than round through f32."
+      },
+      "attrs": { "mode": "constant", "pads": [1, 1] },
+      "inputs": {
+        "data": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [4000000001, 17] } },
+        "constant_value": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [16777217] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "uint32",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [16777217, 4000000001, 17, 16777217] }
+        }
+      }
+    },
+    {
+      "name": "f16_reflect_rank4_inner_aligned",
+      "attrs": { "mode": "reflect", "pads": [0, 0, 2, 2, 0, 0, 2, 2] },
+      "inputs": { "data": { "dtype": "float16", "shape": [1, 16, 4, 8] } },
+      "outputs": { "output": { "dtype": "float16", "shape": [1, 16, 8, 12], "tolerance": 0.003 } }
+    },
+    {
+      "name": "uint8_edge_rank2",
+      "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": {
+        "data": { "dtype": "uint8", "shape": [2, 3], "data": { "kind": "values", "values": [10, 20, 30, 40, 50, 60] } }
+      },
+      "outputs": { "output": { "dtype": "uint8", "shape": [4, 7], "tolerance": 0 } }
+    },
+    {
+      "name": "value_input_reflect_value_ignored",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/pad_test.cc",
+        "test": "PadOpTest.Pad_Reflect_1D",
+        "notes": "constant_value_input gates on all 4 modes; for reflect, map_coord never returns -1 so padValue[0] is never read. Supplies a bogus value (999) to confirm the value input is correctly ignored for non-constant modes."
+      },
+      "attrs": { "mode": "reflect", "pads": [0, 1, 0, 1] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [3, 2],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [999.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "value_input_empty_data_constant_1d",
+      "attrs": { "mode": "constant", "pads": [1, 2] },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [123.0] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [123.0, 123.0, 123.0] }
+        }
+      }
+    },
+    {
+      "name": "value_input_empty_inner_data_2d_constant",
+      "attrs": { "mode": "constant", "pads": [1, 1, 1, 1] },
+      "inputs": {
+        "data": { "dtype": "float32", "shape": [2, 0], "data": { "kind": "values", "values": [] } },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [5.0] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [4, 2],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [5.0, 5.0, 5.0, 5.0, 5.0, 5.0, 5.0, 5.0] }
+        }
+      }
+    },
+    {
+      "name": "value_input_scalar_path_2d_fold_last_row",
+      "attrs": { "mode": "constant", "pads": [1, 0, 1, 0] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [4094, 512],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.023 }
+        },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-7.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4096, 512], "tolerance": 0 } }
+    },
+    {
+      "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": {
+        "data": {
+          "dtype": "float32",
+          "shape": [4, 16, 31],
+          "data": { "kind": "fillFloat32", "sinStep": 0.017, "cosStep": 0.023 }
+        },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-7.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [6, 16, 35], "tolerance": 0 } }
+    },
+    {
+      "name": "rank7_constant_last_axis",
+      "attrs": { "mode": "constant", "pads": [0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 2, 1, 2, 3],
+          "data": { "kind": "linspace", "start": 1.0, "end": 24.0 }
+        },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 1, 2, 1, 2, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "rank8_constant_last_axis",
+      "attrs": { "mode": "constant", "pads": [0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0, 1] },
+      "inputs": {
+        "data": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 2, 1, 2, 2, 3],
+          "data": { "kind": "linspace", "start": 1.0, "end": 48.0 }
+        },
+        "constant_value": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 1, 2, 1, 2, 2, 4], "tolerance": 0 } }
+    }
+  ]
+}