sync 2e7068faf55e

Browse files

Files changed (6) hide show

README.md +73 -0
build/webgpu/bench.json +52 -0
build/webgpu/constant-of-shape.wgsl.jinja +57 -0
build/webgpu/manifest.json +121 -0
build/webgpu/metadata.json +18 -0
build/webgpu/test.json +383 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.ConstantOfShape
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 9
+## Description
+Generates a tensor filled with a constant value and a caller-specified shape. The shape is provided as a 1-D integer tensor; if empty, the output is a scalar. The fill value and its dtype are taken from the `value` attribute, defaulting to `0` of type float32.
+See the [ONNX `ConstantOfShape` spec](https://onnx.ai/onnx/operators/onnx__ConstantOfShape.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | WebGPU storage | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- | --- |
+| `input` | `input` | `S` | `uint32` | `1` | — | Logical int64 1-D tensor whose elements define the output shape; all values must be non-negative and use uint32 WebGPU storage. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `output` | `output` | `T` | derived | — | Output tensor of the shape given by `input`, filled with the constant `value`. | required |
+## Attributes
+Attributes and default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `value` | — | Optional single-element tensor specifying the fill value and output dtype; defaults to 0 of type float32 if omitted. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `S` | `int64` |
+| `T` | `float32`, `float16`, `int32`, `int16`, `uint32`, `int8`, `uint8`, `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
+- [`constant-of-shape.wgsl.jinja`](build/webgpu/constant-of-shape.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.ConstantOfShape", { version: 1 });
+// Explicit destinations request optional results or supply metadata that cannot be inferred.
+const { output } = await kernel({ input: { data: inputData, shape: [2] } }, {
+  outputs: { output: { shape: [2, 3], dtype: "float32" } },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,52 @@

+{
+  "op": "ai.onnx.ConstantOfShape",
+  "cases": [
+    {
+      "name": "mask_16m_f32",
+      "preset": "smoke",
+      "attrs": { "value": -10000 },
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [4096, 4096] } } },
+      "outputs": { "output": { "dtype": "float32", "shape": [4096, 4096] } },
+      "bench": { "primary": true, "metrics": [{ "type": "bandwidth", "value": "4096 * 4096 * 4" }] }
+    },
+    {
+      "name": "1m_f32",
+      "attrs": { "value": -10000 },
+      "inputs": { "input": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1048576] } } },
+      "outputs": { "output": { "dtype": "float32", "shape": [1048576] } }
+    },
+    {
+      "name": "fill_16m_f32_scalar_fallback",
+      "attrs": { "value": -10000 },
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [4097, 4097] } } },
+      "outputs": { "output": { "dtype": "float32", "shape": [4097, 4097] } },
+      "vars": { "dtype": "float32", "count": 16785409 },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype)" }] }
+    },
+    {
+      "name": "mask_16m_f16",
+      "attrs": { "value": -10000 },
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [4096, 4096] } } },
+      "outputs": { "output": { "dtype": "float16", "shape": [4096, 4096] } },
+      "vars": { "dtype": "float16", "count": 16777216 },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype)" }] }
+    },
+    {
+      "name": "fill_134m_f32_2dfold",
+      "attrs": { "value": -10000 },
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [16384, 8192] } } },
+      "outputs": { "output": { "dtype": "float32", "shape": [16384, 8192] } },
+      "vars": { "dtype": "float32", "count": 134217728 },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype)" }] }
+    },
+    {
+      "name": "fill_16m_f16_scalar_fallback",
+      "preset": "stress",
+      "attrs": { "value": -10000 },
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [4097, 4097] } } },
+      "outputs": { "output": { "dtype": "float16", "shape": [4097, 4097] } },
+      "vars": { "dtype": "float16", "count": 16785409 },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype)" }] }
+    }
+  ]
+}

build/webgpu/constant-of-shape.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,57 @@

+{% macro flat_index_2d(name="i", bound="params.count", guardInline=false, note="dispatch-limit") %}
+{% if note == "dispatch-limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
+{% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at the device's per-axis workgroup
+  // limit; gid.y carries the high portion of the output index.
+{% elif note == "vec4-limit" %}
+  // 2D-folded flat vec4 index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills into y).
+{% elif note == "element-limit" %}
+  // 2D-folded flat element index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit.
+{% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension dispatch limit.
+{% endif %}
+{% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+{%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) { return; }
+{%- else %}
+  let {{ name }} = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if ({{ name }} >= {{ bound }}) {
+    return;
+  }
+{%- endif %}
+{% endmacro %}
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+const COUNT: u32 = {{ source.count }}u;
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d("i", "COUNT", note="element-limit") }}
+{% if source.vectorized %}
+{% if source.valueFromUniform %}
+  output[i] = {{ vectorScalar }}({{ scalar }}(params.value));
+{% else %}
+  output[i] = {{ vectorScalar }}({{ scalar }}({{ source.value }}));
+{% endif %}
+{% else %}
+{% if source.valueFromUniform %}
+  output[i] = {{ scalar }}(params.value);
+{% else %}
+  output[i] = {{ scalar }}({{ source.value }});
+{% endif %}
+{% endif %}
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,121 @@

+{
+  "domain": "ai.onnx",
+  "name": "ConstantOfShape",
+  "sinceVersion": 9,
+  "description": "Generates a tensor filled with a constant value and a caller-specified shape. The shape is provided as a 1-D integer tensor; if empty, the output is a scalar. The fill value and its dtype are taken from the `value` attribute, defaulting to `0` of type float32.",
+  "inputs": [
+    {
+      "role": "input",
+      "dtype": "S",
+      "description": "Logical int64 1-D tensor whose elements define the output shape; all values must be non-negative and use uint32 WebGPU storage.",
+      "rank": 1
+    }
+  ],
+  "outputs": [
+    {
+      "role": "output",
+      "dtype": "T",
+      "description": "Output tensor of the shape given by `input`, filled with the constant `value`.",
+      "rank": "dim(shapes.input, 0)"
+    }
+  ],
+  "attributes": {},
+  "attributeDescriptions": {
+    "value": "Optional single-element tensor specifying the fill value and output dtype; defaults to 0 of type float32 if omitted."
+  },
+  "typeConstraints": {
+    "S": ["int64"],
+    "T": ["float32", "float16", "int32", "int16", "uint32", "int8", "uint8", "bool"]
+  },
+  "args": {
+    "input": { "kind": "tensor", "semantic": "input", "role": "shape", "dtype": "uint32", "narrowing": "checked" },
+    "output": { "kind": "tensor", "semantic": "output", "role": "output" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "derive": {
+    "baseContract": "dtypes.S == \"u32\" and ranks.input == 1 and ranks.output >= 0 and ranks.output <= 8 and f16Ok(dtypes.T) and (has(attrs, \"value\") or tensorDtypes.output == \"float32\")",
+    "fillValue": "attrs.value if has(attrs, \"value\") else 0"
+  },
+  "constants": {
+    "scalar": "dtypes.T",
+    "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
+    "outputScalar": "\"vec4<\" ~ dtypes.T ~ \">\" if numel(shapes.output) % 4 == 0 else dtypes.T"
+  },
+  "bindingSets": {
+    "floatFill": [
+      {
+        "name": "output",
+        "arg": "output",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": { "name": "Params", "fields": [{ "name": "value", "type": "f32", "value": "fillValue" }] }
+      }
+    ],
+    "integerFill": [
+      {
+        "name": "output",
+        "arg": "output",
+        "semantic": "output",
+        "buffer": { "type": "storage" },
+        "elementType": "$outputScalar"
+      }
+    ]
+  },
+  "variants": [
+    {
+      "id": "fill_float",
+      "when": ["baseContract", "(dtypes.T == \"f32\" or dtypes.T == \"f16\")"],
+      "constants": { "usesF16": "dtypes.T == \"f16\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConstantOfShape.float",
+          "source": {
+            "shader": "constant-of-shape.wgsl.jinja",
+            "inputs": {
+              "count": "numel(shapes.output) / 4 if numel(shapes.output) % 4 == 0 else numel(shapes.output)",
+              "valueFromUniform": true,
+              "vectorized": "numel(shapes.output) % 4 == 0"
+            }
+          },
+          "bindings": "floatFill",
+          "dispatch": {
+            "threads": "numel(shapes.output) / 4 if numel(shapes.output) % 4 == 0 else numel(shapes.output)",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        }
+      ]
+    },
+    {
+      "id": "fill_integer",
+      "when": ["baseContract", "dtypes.T != \"f32\" and dtypes.T != \"f16\""],
+      "constants": { "usesF16": false },
+      "passes": [
+        {
+          "id": "main",
+          "name": "ConstantOfShape.integer",
+          "source": {
+            "shader": "constant-of-shape.wgsl.jinja",
+            "inputs": {
+              "count": "numel(shapes.output) / 4 if numel(shapes.output) % 4 == 0 else numel(shapes.output)",
+              "value": "fillValue",
+              "valueFromUniform": false,
+              "vectorized": "numel(shapes.output) % 4 == 0"
+            }
+          },
+          "bindings": "integerFill",
+          "dispatch": {
+            "threads": "numel(shapes.output) / 4 if numel(shapes.output) % 4 == 0 else numel(shapes.output)",
+            "workgroupSize": "tunables.WORKGROUP_SIZE"
+          }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,18 @@

+{
+  "name": "ai.onnx.ConstantOfShape",
+  "id": "_ai_onnx_constantofshape_webgpu_8abe390",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "hw5exM30DtUDB3GcNZRZZLLlss8j08uRGSj+WowzDRI=",
+      "constant-of-shape.wgsl.jinja": "vIBoi4zlY5R5eGCNEUQdtotvQnPKLMokdYPvBLr39eg=",
+      "manifest.json": "bvhlqmL0sQ4VC3S3twA7PQVbUCmyyQ3H5R6IWmYwEWg=",
+      "test.json": "8Y206ANaAdGdYOYNjOOQuuwZ31uxxZx9ZPMuhFzMNWc="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.ConstantOfShape" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,383 @@

+{
+  "op": "ai.onnx.ConstantOfShape",
+  "usageCase": "zero_default",
+  "cases": [
+    {
+      "name": "zero_default",
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 3] } } },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3] } }
+    },
+    {
+      "name": "f32_vec4_fill_4096",
+      "provenance": {
+        "notes": "Compact sibling for the f32 ConstantOfShape fill benchmark; preserves a vec4-aligned flat fill without benchmark-scale output."
+      },
+      "attrs": { "value": -10000 },
+      "inputs": { "input": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [4096] } } },
+      "outputs": { "output": { "dtype": "float32", "shape": [4096], "tolerance": 0 } }
+    },
+    {
+      "name": "f32_subnormal_fill_value",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/constant_of_shape_test.cc",
+        "test": "ConstantOfShape.Float_Ones",
+        "notes": "Valid float32 fill value below the normal range; the uniform-to-storage fill path must preserve finite subnormal payloads."
+      },
+      "attrs": { "value": 1e-40 },
+      "inputs": { "input": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [4] } } },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [1e-40, 1e-40, 1e-40, 1e-40] }
+        }
+      }
+    },
+    {
+      "name": "f32_positive_infinity_fill",
+      "attrs": { "value": "Infinity" },
+      "inputs": { "input": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [4] } } },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": ["Infinity", "Infinity", "Infinity", "Infinity"] }
+        }
+      }
+    },
+    {
+      "name": "f32_negative_infinity_fill",
+      "attrs": { "value": "-Infinity" },
+      "inputs": { "input": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [3] } } },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [3],
+          "data": { "kind": "values", "values": ["-Infinity", "-Infinity", "-Infinity"] }
+        }
+      }
+    },
+    {
+      "name": "f16_fill",
+      "attrs": { "value": 1.5 },
+      "inputs": { "input": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [1, 2, 4] } } },
+      "outputs": { "output": { "dtype": "float16", "shape": [1, 2, 4] } },
+      "tolerance": 0.001
+    },
+    {
+      "name": "scalar_output_empty_shape",
+      "attrs": { "value": -4.25 },
+      "inputs": { "input": { "dtype": "uint32", "shape": [0], "data": { "kind": "values", "values": [] } } },
+      "outputs": { "output": { "dtype": "float32", "shape": [], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank7_output_fill",
+      "attrs": { "value": -2.75 },
+      "inputs": {
+        "input": { "dtype": "uint32", "shape": [7], "data": { "kind": "values", "values": [1, 2, 1, 3, 1, 2, 2] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 1, 3, 1, 2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "zero_sized_output_fill",
+      "attrs": { "value": 6.5 },
+      "inputs": { "input": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [2, 0, 3] } } },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 0, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_int32_zeros_10x6",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/constant_of_shape_test.cc",
+        "test": "ConstantOfShape.Int32_Zeros"
+      },
+      "attrs": { "value": 0 },
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [10, 6] } } },
+      "outputs": { "output": { "dtype": "int32", "shape": [10, 6], "tolerance": 0 } }
+    },
+    {
+      "name": "int32_negative_fill_exact",
+      "attrs": { "value": -17 },
+      "inputs": { "input": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [2, 1, 3] } } },
+      "outputs": { "output": { "dtype": "int32", "shape": [2, 1, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "uint32_fill_exact_above_float24",
+      "attrs": { "value": 16777217 },
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 3] } } },
+      "outputs": { "output": { "dtype": "uint32", "shape": [2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_float_ones_4x3x2",
+      "attrs": { "value": 1 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/constant_of_shape_test.cc",
+        "test": "ConstantOfShape.Float_Ones"
+      },
+      "inputs": { "input": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [4, 3, 2] } } },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 3, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_default_value_float_zeros_2x6",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/constant_of_shape_test.cc",
+        "test": "ConstantOfShape.DefaultValue"
+      },
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 6] } } },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 6], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_type_test_float16_bits_value",
+      "attrs": { "value": 2.980232238769531e-7 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/constant_of_shape_test.cc",
+        "test": "ConstantOfShape.TypeTests",
+        "notes": "FLOAT16 typed value from raw bits 5, represented as its numeric half value."
+      },
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 6] } } },
+      "outputs": { "output": { "dtype": "float16", "shape": [2, 6], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_type_test_uint32_2x6",
+      "attrs": { "value": 32 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/constant_of_shape_test.cc",
+        "test": "ConstantOfShape.TypeTests",
+        "notes": "UINT32 typed value."
+      },
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 6] } } },
+      "outputs": { "output": { "dtype": "uint32", "shape": [2, 6], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_type_test_int32_exact_above_float24",
+      "attrs": { "value": 20000001 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/constant_of_shape_test.cc",
+        "test": "ConstantOfShape.TypeTests",
+        "notes": "INT32 typed value variant extended above f32's exact integer range."
+      },
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 3] } } },
+      "outputs": { "output": { "dtype": "int32", "shape": [2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_type_test_int16_2x6",
+      "attrs": { "value": 16 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/constant_of_shape_test.cc",
+        "test": "ConstantOfShape.TypeTests",
+        "notes": "INT16 typed value from ORT's ConstantOfShape TypeTests."
+      },
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 6] } } },
+      "outputs": { "output": { "dtype": "int16", "shape": [2, 6], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_type_test_int16_min_fill",
+      "attrs": { "value": -32768 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/constant_of_shape_test.cc",
+        "test": "ConstantOfShape.TypeTests",
+        "notes": "Signed int16 sentinel extension of ORT's typed ConstantOfShape TypeTests."
+      },
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 3] } } },
+      "outputs": {
+        "output": {
+          "dtype": "int16",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [-32768, -32768, -32768, -32768, -32768, -32768] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "ort_type_test_uint32_large_exact",
+      "attrs": { "value": 4000000001 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/constant_of_shape_test.cc",
+        "test": "ConstantOfShape.TypeTests",
+        "notes": "UINT32 typed value variant extended above f32's exact integer range."
+      },
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 3] } } },
+      "outputs": { "output": { "dtype": "uint32", "shape": [2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_type_test_int8_min_fill",
+      "attrs": { "value": -128 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/constant_of_shape_test.cc",
+        "test": "ConstantOfShape.TypeTests",
+        "notes": "Exercises the same typed TensorProto value path as ORT's int8 TypeTests with a signed edge value."
+      },
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 6] } } },
+      "outputs": { "output": { "dtype": "int8", "shape": [2, 6], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_type_test_uint8_high_fill",
+      "attrs": { "value": 250 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/constant_of_shape_test.cc",
+        "test": "ConstantOfShape.TypeTests",
+        "notes": "Exercises the same typed TensorProto value path as ORT's uint8 TypeTests with a high logical byte value."
+      },
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 6] } } },
+      "outputs": { "output": { "dtype": "uint8", "shape": [2, 6], "tolerance": 0 } }
+    },
+    {
+      "name": "bool_true_fill",
+      "attrs": { "value": 1 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/constant_of_shape_test.cc",
+        "test": "ConstantOfShape.TypeTests",
+        "notes": "ORT source notes bool was skipped only because its C++ test harness cannot use vector<bool> as a contiguous buffer."
+      },
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 3] } } },
+      "outputs": {
+        "output": {
+          "dtype": "bool",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [1, 1, 1, 1, 1, 1] },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "onnx_backend_constantofshape_float_ones",
+      "inputs": { "input": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [4, 3, 2] } } },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 3, 2] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_constantofshape_float_ones",
+        "notes": "ONNX int64 metadata/index tensors use framework int32/uint32 slots where representable."
+      }
+    },
+    {
+      "name": "onnx_backend_constantofshape_int_shape_zero",
+      "inputs": { "input": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [0] } } },
+      "outputs": { "output": { "dtype": "int32", "shape": [0] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_constantofshape_int_shape_zero",
+        "notes": "ONNX int64 metadata/index tensors use framework int32/uint32 slots where representable. ONNX tensor-valued value attribute represented as scalar framework value attr."
+      },
+      "attrs": { "value": 0 }
+    },
+    {
+      "name": "onnx_backend_constantofshape_int_zeros",
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [10, 6] } } },
+      "outputs": { "output": { "dtype": "int32", "shape": [10, 6] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_constantofshape_int_zeros",
+        "notes": "ONNX int64 metadata/index tensors use framework int32/uint32 slots where representable. ONNX tensor-valued value attribute represented as scalar framework value attr."
+      },
+      "attrs": { "value": 0 }
+    },
+    {
+      "name": "rank8_output_scalar_fill",
+      "attrs": { "value": 3.5 },
+      "provenance": {
+        "source": "onnx spec: ConstantOfShape output rank = len(shape input), unbounded",
+        "notes": "Fills a rank-8 tensor with the scalar float32 value 3.5."
+      },
+      "inputs": {
+        "input": { "dtype": "uint32", "shape": [8], "data": { "kind": "values", "values": [2, 1, 2, 1, 2, 1, 2, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 1, 2, 1, 2, 1, 2, 1], "tolerance": 0 } }
+    },
+    {
+      "name": "rank8_output_oddnumel_fill",
+      "attrs": { "value": -1.25 },
+      "provenance": {
+        "source": "onnx spec: ConstantOfShape arbitrary output rank",
+        "notes": "Rank-8 float32 fill with numel 18 (%4 != 0 -> would target the fill_float SCALAR fallback). Distinct from the vec4 case so the fix is pinned on BOTH the vec4 and the scalar fill paths. ORT CPU computes it; GPU throws because every variant gates ranks.output <= 7."
+      },
+      "inputs": {
+        "input": { "dtype": "uint32", "shape": [8], "data": { "kind": "values", "values": [2, 1, 3, 1, 1, 3, 1, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 1, 3, 1, 1, 3, 1, 1], "tolerance": 0 } }
+    },
+    {
+      "name": "f16_oddnumel_scalar_fill",
+      "attrs": { "value": -2.5 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/constant_of_shape_test.cc (TypeTests float16)",
+        "notes": "A 15-element float16 output exercises the scalar fill path because its element count is not divisible by four. The value is exactly representable in float16."
+      },
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 5] } } },
+      "outputs": { "output": { "dtype": "float16", "shape": [3, 5], "tolerance": 0.001 } }
+    },
+    {
+      "name": "u32_oddnumel_scalar_fill",
+      "attrs": { "value": 7 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/constant_of_shape_test.cc (TypeTests uint32)",
+        "notes": "Clean coverage: uint32 output with numel 15 (%4 != 0) selects the scalar branch of fill_integer. Existing uint32 cases are vec4-aligned, so this pins the scalar integer fill path."
+      },
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 5] } } },
+      "outputs": { "output": { "dtype": "uint32", "shape": [3, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "int32_vec4_fill_16m_2dfold_correctness",
+      "attrs": { "value": 123456789 },
+      "provenance": {
+        "source": "onnx spec: ConstantOfShape flat memset; 2D dispatch-fold guard",
+        "notes": "Vec4 int32 fill with numel 67108864 selects fill_integer's vectorized branch. numel/4 = 16777216 > 65535*256, so gid.y carries the high bits. Verifies the fold index+guard writes every element including the gid.y>0 region. Every element must equal 123456789."
+      },
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [8192, 8192] } } },
+      "outputs": {
+        "output": {
+          "dtype": "int32",
+          "shape": [8192, 8192],
+          "data": { "kind": "constant", "value": 123456789 },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "i32_vec4_min_int32_fill_exact",
+      "attrs": { "value": -2147483648 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/constant_of_shape_test.cc (TypeTests int32) extended to INT32_MIN",
+        "notes": "Numel 8 selects fill_integer's vectorized branch. The full-width signed sentinel verifies the value is specialized as i32 with no f32 round-trip. Every element must equal -2147483648."
+      },
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 4] } } },
+      "outputs": {
+        "output": {
+          "dtype": "int32",
+          "shape": [2, 4],
+          "data": {
+            "kind": "values",
+            "values": [-2147483648, -2147483648, -2147483648, -2147483648, -2147483648, -2147483648, -2147483648, -2147483648]
+          },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "u32_vec4_max_uint32_fill_exact",
+      "attrs": { "value": 4294967295 },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/generator/constant_of_shape_test.cc (TypeTests uint32) extended to UINT32_MAX",
+        "notes": "Numel 8 selects fill_integer's vectorized branch. The full-width unsigned sentinel verifies the value is specialized as u32 with no f32 round-trip (which would wrap 4294967295 to 0). Every element must equal 4294967295."
+      },
+      "inputs": { "input": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 4] } } },
+      "outputs": {
+        "output": {
+          "dtype": "uint32",
+          "shape": [2, 4],
+          "data": {
+            "kind": "values",
+            "values": [4294967295, 4294967295, 4294967295, 4294967295, 4294967295, 4294967295, 4294967295, 4294967295]
+          },
+          "tolerance": 0
+        }
+      }
+    },
+    {
+      "name": "i32_vec4_dim1_first_zero_sized_output",
+      "attrs": { "value": -17 },
+      "provenance": {
+        "source": "onnx spec: ConstantOfShape with a zero extent -> empty output; array<T,0> compile-crash class",
+        "notes": "shape [0,5,3] -> numel 0, dtypes.T==i32. Verifies the i32 vec4/scalar fill path does not crash on a leading-zero empty output. Existing zero_sized_output_fill covers f32 with a middle zero."
+      },
+      "inputs": { "input": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [0, 5, 3] } } },
+      "outputs": { "output": { "dtype": "int32", "shape": [0, 5, 3], "tolerance": 0 } }
+    }
+  ]
+}