sync 2e7068faf55e

Browse files

Files changed (7) hide show

README.md +70 -0
build/webgpu/bench.json +42 -0
build/webgpu/expand-vec4.wgsl.jinja +88 -0
build/webgpu/expand.wgsl.jinja +50 -0
build/webgpu/manifest.json +128 -0
build/webgpu/metadata.json +19 -0
build/webgpu/test.json +607 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.Expand
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 13
+## Description
+Broadcasts `input` to a given `shape` following NumPy-style rules: dimensions are right-aligned, and corresponding dimensions must be equal or one of them must be 1. The output shape may differ from `shape` when a requested dimension is 1 or `shape` has fewer dimensions than `input`.
+See the [ONNX `Expand` spec](https://onnx.ai/onnx/operators/onnx__Expand.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | WebGPU storage | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- | --- |
+| `input` | `input` | `T` | runtime-selected; narrow integers and bool use 32-bit slots | — | — | Input tensor to broadcast. | required |
+| `shape` | `shape` | `S` | `uint32` | `1` | — | Logical int64 1-D tensor specifying the non-negative target shape; WebGPU stores it as uint32. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `output` | `output` | `T` | derived | — | Output tensor with the broadcasted shape. | required |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16`, `uint32`, `int32`, `int16`, `uint8`, `int8`, `bool` |
+| `S` | `int64` |
+## 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
+- [`expand-vec4.wgsl.jinja`](build/webgpu/expand-vec4.wgsl.jinja)
+- [`expand.wgsl.jinja`](build/webgpu/expand.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.Expand", { version: 1 });
+// Explicit destinations request optional results or supply metadata that cannot be inferred.
+const { output } = await kernel({
+  input: { data: inputData, shape: [1, 4] },
+  shape: { data: shapeData, shape: [2] },
+}, {
+  outputs: { output: { shape: [2, 4], dtype: "int16" } },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,42 @@

+{
+  "op": "ai.onnx.Expand",
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "expand-f32-1m",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "count": 1048576 },
+      "inputs": {
+        "input": { "shape": [1, 1024], "dtype": "float32", "dist": "normal", "seed": 860, "scale": 1 },
+        "shape": { "shape": [2], "dtype": "uint32", "data": { "kind": "values", "values": [1024, 1024] } }
+      },
+      "outputs": { "output": { "shape": [1024, 1024], "dtype": "float32" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }]
+      }
+    },
+    {
+      "name": "expand-u32-bool-like-1m",
+      "preset": "edge",
+      "vars": { "dtype": "uint32", "count": 1048576 },
+      "inputs": {
+        "input": { "shape": [1, 1024], "dtype": "uint32", "dist": "linearMod", "mod": 2 },
+        "shape": { "shape": [2], "dtype": "uint32", "data": { "kind": "values", "values": [1024, 1024] } }
+      },
+      "outputs": { "output": { "shape": [1024, 1024], "dtype": "uint32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "expand-f32-scalar-fallback-1m",
+      "preset": "stress",
+      "vars": { "dtype": "float32", "count": 1049600 },
+      "inputs": {
+        "input": { "shape": [1, 1025], "dtype": "float32", "dist": "normal", "seed": 42, "scale": 1 },
+        "shape": { "shape": [2], "dtype": "uint32", "data": { "kind": "values", "values": [1024, 1025] } }
+      },
+      "outputs": { "output": { "shape": [1024, 1025], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    }
+  ]
+}

build/webgpu/expand-vec4.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,88 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+{% set x_numel = namespace(value=1) %}
+{% for d in source.xShape %}{% set x_numel.value = x_numel.value * d %}{% endfor %}
+{% set x_same = namespace(value=(source.xRank == source.yRank)) %}
+{% if x_same.value %}
+{% for axis in range(source.yRank) %}
+{% if source.xShape[axis] != source.yShape[axis] %}{% set x_same.value = false %}{% endif %}
+{% endfor %}
+{% endif %}
+{% set x_inner = source.xShape[source.xRank - 1] if source.xRank >= 1 else 1 %}
+{% if x_numel.value == 1 %}{% set x_mode = "scalar" %}
+{% elif x_inner == 1 %}{% set x_mode = "splat" %}
+{% else %}{% set x_mode = "contig" %}{% endif %}
+{% set xUsesIndex = namespace(value=false) %}
+{% for axis in range(source.yRank) %}
+{% set x_axis = axis - (source.yRank - source.xRank) %}
+{% if x_axis >= 0 and source.xShape[x_axis] != 1 %}{% set xUsesIndex.value = true %}{% endif %}
+{% endfor %}
+{% if x_mode == "scalar" %}
+// Each invocation splats the single input value across four consecutive outputs.
+{% elif x_mode == "splat" %}
+// Each invocation broadcasts one innermost-axis value across four outputs;
+// x_offset applies any outer-axis broadcasting once for the group.
+{% else %}
+// Each invocation copies four contiguous input values. x_offset applies any
+// outer-axis broadcasting once for the group.
+{% endif %}
+{% if x_mode != "scalar" %}
+fn x_offset(out_index: u32) -> u32 {
+{% if x_same.value %}
+  return out_index;
+{% else %}
+{% if xUsesIndex.value %}
+  var rem = out_index;
+{% endif %}
+  var offset = 0u;
+  // Empty output axes can produce a zero trailing stride. Their dispatch is
+  // empty, but the module must still compile, so zero strides use a harmless
+  // divisor of one instead of emitting a constant division by zero.
+{% for axis in range(source.yRank) %}
+{% set y_stride = namespace(value=1) %}
+{% for j in range(axis + 1, source.yRank) %}
+{% set y_stride.value = y_stride.value * source.yShape[j] %}
+{% endfor %}
+{% set y_div = y_stride.value if y_stride.value != 0 else 1 %}
+{% set x_axis = axis - (source.yRank - source.xRank) %}
+{% if x_axis >= 0 and source.xShape[x_axis] != 1 %}
+  let coord{{ axis }} = rem / {{ y_div }}u;
+  rem = rem % {{ y_div }}u;
+{% set x_stride = namespace(value=1) %}
+{% for j in range(x_axis + 1, source.xRank) %}
+{% set x_stride.value = x_stride.value * source.xShape[j] %}
+{% endfor %}
+  offset = offset + coord{{ axis }} * {{ x_stride.value }}u;
+{% elif xUsesIndex.value %}
+  rem = rem % {{ y_div }}u;
+{% endif %}
+{% endfor %}
+  return offset;
+{% endif %}
+}
+{% endif %}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+  // 2D-folded flat vec4 index: gid.y carries the high bits past the
+  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills into y).
+  let i4 = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if (i4 >= params.count) {
+    return;
+  }
+{% if x_mode != "scalar" %}
+  let base = i4 * 4u;
+{% endif %}
+{% if x_mode == "scalar" %}
+  y[i4] = {{ vectorScalar }}(x[0]);
+{% elif x_mode == "splat" %}
+  y[i4] = {{ vectorScalar }}(x[x_offset(base)]);
+{% else %}
+  let xo = x_offset(base);
+  y[i4] = {{ vectorScalar }}(x[xo], x[xo + 1u], x[xo + 2u], x[xo + 3u]);
+{% endif %}
+}

build/webgpu/expand.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,50 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+{% set xUsesIndex = namespace(value=false) %}
+{% for axis in range(source.yRank) %}
+{% set x_axis = axis - (source.yRank - source.xRank) %}
+{% if x_axis >= 0 and source.xShape[x_axis] != 1 %}{% set xUsesIndex.value = true %}{% endif %}
+{% endfor %}
+fn x_offset({% if xUsesIndex.value %}out_index: u32{% endif %}) -> u32 {
+{% if xUsesIndex.value %}
+  var rem = out_index;
+{% endif %}
+  var offset = 0u;
+  // Empty output axes can produce a zero trailing stride. Their dispatch is
+  // empty, but the module must still compile, so zero strides use a harmless
+  // divisor of one instead of emitting a constant division by zero.
+{% for axis in range(source.yRank) %}
+{% set y_stride = namespace(value=1) %}
+{% for j in range(axis + 1, source.yRank) %}
+{% set y_stride.value = y_stride.value * source.yShape[j] %}
+{% endfor %}
+{% set y_div = y_stride.value if y_stride.value != 0 else 1 %}
+{% set x_axis = axis - (source.yRank - source.xRank) %}
+{% if x_axis >= 0 and source.xShape[x_axis] != 1 %}
+  let coord{{ axis }} = rem / {{ y_div }}u;
+  rem = rem % {{ y_div }}u;
+{% set x_stride = namespace(value=1) %}
+{% for j in range(x_axis + 1, source.xRank) %}
+{% set x_stride.value = x_stride.value * source.xShape[j] %}
+{% endfor %}
+  offset = offset + coord{{ axis }} * {{ x_stride.value }}u;
+{% elif xUsesIndex.value %}
+  rem = rem % {{ y_div }}u;
+{% endif %}
+{% 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 dispatch limit (outputs > 16.7M elements).
+  let i = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if (i >= params.count) {
+    return;
+  }
+  y[i] = x[x_offset({% if xUsesIndex.value %}i{% endif %})];
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,128 @@

+{
+  "domain": "ai.onnx",
+  "name": "Expand",
+  "sinceVersion": 13,
+  "description": "Broadcasts `input` to a given `shape` following NumPy-style rules: dimensions are right-aligned, and corresponding dimensions must be equal or one of them must be 1. The output shape may differ from `shape` when a requested dimension is 1 or `shape` has fewer dimensions than `input`.",
+  "inputs": [
+    { "role": "input", "dtype": "T", "description": "Input tensor to broadcast." },
+    {
+      "role": "shape",
+      "dtype": "S",
+      "description": "Logical int64 1-D tensor specifying the non-negative target shape; WebGPU stores it as uint32.",
+      "rank": 1
+    }
+  ],
+  "outputs": [
+    {
+      "role": "output",
+      "dtype": "T",
+      "description": "Output tensor with the broadcasted shape.",
+      "rank": "dim(shapes.shape, 0)"
+    }
+  ],
+  "typeConstraints": {
+    "T": ["float32", "float16", "uint32", "int32", "int16", "uint8", "int8", "bool"],
+    "S": ["int64"]
+  },
+  "args": {
+    "input": { "kind": "tensor", "semantic": "input", "role": "input" },
+    "shape": { "kind": "tensor", "semantic": "shape", "role": "shape", "dtype": "uint32", "narrowing": "checked" },
+    "output": { "kind": "tensor", "semantic": "output", "role": "output" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
+  "variants": [
+    {
+      "id": "broadcast_vec4",
+      "priority": 10,
+      "when": ["ranks.output >= 1", "ranks.input <= ranks.output", "f16Ok(dtypes.T)", "dim(shapes.output, -1) % 4 == 0"],
+      "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Expand.vec4",
+          "source": {
+            "shader": "expand-vec4.wgsl.jinja",
+            "inputs": {
+              "xShape": "shapes.input",
+              "yShape": "shapes.output",
+              "xRank": "ranks.input",
+              "yRank": "ranks.output"
+            }
+          },
+          "bindings": [
+            {
+              "name": "x",
+              "arg": "input",
+              "semantic": "input",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar"
+            },
+            {
+              "name": "y",
+              "arg": "output",
+              "semantic": "output",
+              "buffer": { "type": "storage" },
+              "elementType": "$vectorScalar"
+            },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.output) / 4" }]
+              }
+            }
+          ],
+          "dispatch": { "threads": "numel(shapes.output) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "broadcast",
+      "when": ["ranks.input <= ranks.output", "f16Ok(dtypes.T)"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "Expand",
+          "source": {
+            "shader": "expand.wgsl.jinja",
+            "inputs": {
+              "xShape": "shapes.input",
+              "yShape": "shapes.output",
+              "xRank": "ranks.input",
+              "yRank": "ranks.output"
+            }
+          },
+          "bindings": [
+            {
+              "name": "x",
+              "arg": "input",
+              "semantic": "input",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar"
+            },
+            {
+              "name": "y",
+              "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)" }]
+              }
+            }
+          ],
+          "dispatch": { "threads": "numel(shapes.output)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

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

+{
+  "name": "ai.onnx.Expand",
+  "id": "_ai_onnx_expand_webgpu_54cd655",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "RLidqSlxqJDOxs+FeBKGm1VWNsXk2nLH7N+hyCych0Q=",
+      "expand-vec4.wgsl.jinja": "aiqKbsXQ+vUgxM55kKrOrVbpcLgvJzXSenpuGx+/C/w=",
+      "expand.wgsl.jinja": "LurmTZjbk5/Jtet50CvC5vPLooK/k4cPHmlEWHi3k64=",
+      "manifest.json": "GHAUTVFssEJLTI/FFhC0SV610+2xmx65imK6FCmKRk0=",
+      "test.json": "T6kUaQLTGfJvF6V4kE+IwKWM2zmZu8QuPHyo+iy+HJs="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Expand" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,607 @@

+{
+  "op": "ai.onnx.Expand",
+  "cases": [
+    {
+      "name": "int16_broadcast_boundaries",
+      "inputs": {
+        "input": { "dtype": "int16", "shape": [1, 4], "data": { "kind": "values", "values": [-32768, -1, 0, 32767] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 4] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "int16",
+          "shape": [2, 4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [-32768, -1, 0, 32767, -32768, -1, 0, 32767] }
+        }
+      }
+    },
+    {
+      "name": "dispatch_cliff_scalar_splat",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [7.0] } },
+        "shape": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [16776963] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [16776963], "tolerance": 0 } }
+    },
+    {
+      "name": "scalar_to_3x3",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [2.5] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 3] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_float_scalar_to_3x3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_3x3"
+      },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 3] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "row_vector_to_3x3",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1, 3], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 3] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_float_vector_to_rows",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_3x1"
+      },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "col_vector_to_3x3",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [3, 1], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 3] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_float_col_vector_to_cols",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_1x3"
+      },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [3, 1], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [1, 3] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank3_broadcast",
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 3, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.13, "cosStep": 0.29 }
+        },
+        "shape": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [2, 3, 4] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "uint32_bool_like_broadcast",
+      "inputs": {
+        "input": { "dtype": "uint32", "shape": [1, 3, 1], "data": { "kind": "values", "values": [1, 0, 2] } },
+        "shape": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [2, 3, 4] } }
+      },
+      "outputs": { "output": { "dtype": "uint32", "shape": [2, 3, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "rank5_broadcast",
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 1, 2, 1, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.11, "cosStep": 0.23 }
+        },
+        "shape": { "dtype": "uint32", "shape": [5], "data": { "kind": "values", "values": [2, 3, 2, 4, 2] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 2, 4, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank7_broadcast",
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 3, 1, 1, 2],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.19 }
+        },
+        "shape": { "dtype": "uint32", "shape": [7], "data": { "kind": "values", "values": [2, 2, 4, 3, 2, 3, 2] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2, 4, 3, 2, 3, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "true_scalar_to_rank3",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-2.5] } },
+        "shape": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [2, 2, 3] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2, 3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "scalar_to_scalar_empty_shape",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [6.5] } },
+        "shape": { "dtype": "uint32", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "broadcast_to_zero_width_rows",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_3x1",
+        "notes": "Valid empty-output broadcast: input shape [1,0] expands to [3,0]. No elements are written, but generated stride math must still compile."
+      },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1, 0], "data": { "kind": "values", "values": [] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 0] } }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [3, 0], "tolerance": 0, "data": { "kind": "values", "values": [] } }
+      }
+    },
+    {
+      "name": "broadcast_to_zero_middle_dim_scalar",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_3x1x3x1_int64",
+        "notes": "Scalar-template companion for zero-sized broadcast output: the final dimension is not vec4-aligned, so the generic Expand shader must handle zero strides."
+      },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [1, 0, 5], "data": { "kind": "values", "values": [] } },
+        "shape": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [3, 0, 5] } }
+      },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [3, 0, 5], "tolerance": 0, "data": { "kind": "values", "values": [] } }
+      }
+    },
+    {
+      "name": "ort_int32_scalar_to_3x3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_3x3_int32"
+      },
+      "inputs": {
+        "input": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 3] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int32_vector_to_rows",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_3x1_int32"
+      },
+      "inputs": {
+        "input": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [1, 2, 3] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 1] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int32_rank4_singletons",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_3x1x3x1_int64",
+        "notes": "Same rank-4 singleton broadcast pattern adapted to int32 storage."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "int32",
+          "shape": [1, 3, 1, 3],
+          "data": { "kind": "values", "values": [1, 2, 3, 4, 5, 6, 7, 8, 9] }
+        },
+        "shape": { "dtype": "uint32", "shape": [4], "data": { "kind": "values", "values": [3, 3, 3, 3] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [3, 3, 3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_f16_scalar_to_3x3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_3x3_float16"
+      },
+      "inputs": {
+        "input": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [1.0] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 3] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [3, 3], "tolerance": 0.001 } }
+    },
+    {
+      "name": "ort_f16_vector_to_rows",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_3x1_float16"
+      },
+      "inputs": {
+        "input": { "dtype": "float16", "shape": [3], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 1] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [3, 3], "tolerance": 0.001 } }
+    },
+    {
+      "name": "ort_f16_col_vector_to_cols",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_1x3_float16"
+      },
+      "inputs": {
+        "input": { "dtype": "float16", "shape": [3, 1], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [1, 3] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [3, 3], "tolerance": 0.001 } }
+    },
+    {
+      "name": "ort_uint32_bool_scalar_to_3x3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_3x3_bool",
+        "notes": "Adapts ORT's bool tensor to logical uint32 values because this framework stores bool-like tensors as u32 slots."
+      },
+      "inputs": {
+        "input": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 3] } }
+      },
+      "outputs": { "output": { "dtype": "uint32", "shape": [3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_uint32_bool_col_to_four_cols",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_1x4_bool",
+        "notes": "Adapts ORT's bool tensor to logical uint32 values because this framework stores bool-like tensors as u32 slots."
+      },
+      "inputs": {
+        "input": { "dtype": "uint32", "shape": [3, 1], "data": { "kind": "values", "values": [0, 1, 0] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [1, 4] } }
+      },
+      "outputs": { "output": { "dtype": "uint32", "shape": [3, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_shape_dim_one_does_not_shrink_rank5",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_2x2x1x2x1_float"
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 2, 1, 2, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0] }
+        },
+        "shape": { "dtype": "uint32", "shape": [5], "data": { "kind": "values", "values": [1, 2, 2, 2, 2] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2, 2, 2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_shape_dim_one_does_not_shrink_middle",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_3x1x8_float"
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [3, 2, 1],
+          "data": { "kind": "values", "values": [1.0, 2.0, 3.0, 4.0, 5.0, 6.0] }
+        },
+        "shape": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [3, 1, 8] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 2, 8], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_int32_col_vector_to_cols",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_1x3_int32"
+      },
+      "inputs": {
+        "input": { "dtype": "int32", "shape": [3, 1], "data": { "kind": "values", "values": [1, 2, 3] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [1, 3] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_uint32_bool_col_to_cols",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_1x3_bool",
+        "notes": "Adapts ORT's bool tensor to logical uint32 values because this framework stores bool-like tensors as u32 slots."
+      },
+      "inputs": {
+        "input": { "dtype": "uint32", "shape": [3, 1], "data": { "kind": "values", "values": [0, 1, 0] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [1, 3] } }
+      },
+      "outputs": { "output": { "dtype": "uint32", "shape": [3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_uint32_bool_row_to_rows",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_4x1_bool",
+        "notes": "Adapts ORT's bool tensor to logical uint32 values because this framework stores bool-like tensors as u32 slots."
+      },
+      "inputs": {
+        "input": { "dtype": "uint32", "shape": [1, 4], "data": { "kind": "values", "values": [0, 1, 0, 0] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [4, 1] } }
+      },
+      "outputs": { "output": { "dtype": "uint32", "shape": [4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_scalar_float_empty_shape",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_scalar_float"
+      },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [3.0] } },
+        "shape": { "dtype": "uint32", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "onnx_backend_expand_dim_changed",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [3, 1], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } },
+        "shape": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [2, 1, 6] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 6] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_expand_dim_changed",
+        "notes": "ONNX int64 metadata/index tensors use framework int32/uint32 slots where representable."
+      }
+    },
+    {
+      "name": "onnx_backend_expand_dim_unchanged",
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [3, 1], "data": { "kind": "values", "values": [1.0, 2.0, 3.0] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 4] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 4] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_expand_dim_unchanged",
+        "notes": "ONNX int64 metadata/index tensors use framework int32/uint32 slots where representable."
+      }
+    },
+    {
+      "name": "ort_scalar_float_to_scalar",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_scalar_float",
+        "notes": "ORT shape is int64; this framework stores the empty shape tensor in uint32 slots."
+      },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [3.0] } },
+        "shape": { "dtype": "uint32", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_scalar_int32_to_rank3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_scalar_int32",
+        "notes": "ORT shape is int64; this framework stores representable shape values in uint32 slots."
+      },
+      "inputs": {
+        "input": { "dtype": "int32", "shape": [], "data": { "kind": "values", "values": [9] } },
+        "shape": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [2, 3, 4] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [2, 3, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_bool_pattern_uint32_row_expand",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_3x1_bool",
+        "notes": "ORT uses bool payloads; this framework represents logical bool-like values as uint32 slots."
+      },
+      "inputs": {
+        "input": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [0, 1, 0] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 1] } }
+      },
+      "outputs": { "output": { "dtype": "uint32", "shape": [3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int8_vector_to_rows_edge_values",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_3x1_int32",
+        "notes": "Same row-broadcast shape as ORT's int32 case, using ONNX-valid signed byte edge values."
+      },
+      "inputs": {
+        "input": { "dtype": "int8", "shape": [3], "data": { "kind": "values", "values": [-128, -1, 127] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 1] } }
+      },
+      "outputs": { "output": { "dtype": "int8", "shape": [3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_uint8_col_vector_to_cols_edge_values",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_1x3_int32",
+        "notes": "Same column-broadcast shape as ORT's int32 case, using ONNX-valid unsigned byte edge values."
+      },
+      "inputs": {
+        "input": { "dtype": "uint8", "shape": [3, 1], "data": { "kind": "values", "values": [0, 128, 255] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [1, 3] } }
+      },
+      "outputs": { "output": { "dtype": "uint8", "shape": [3, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_bool_scalar_to_3x3",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_3x3_bool"
+      },
+      "inputs": {
+        "input": { "dtype": "bool", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 3] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "bool",
+          "shape": [3, 3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [1, 1, 1, 1, 1, 1, 1, 1, 1] }
+        }
+      }
+    },
+    {
+      "name": "ort_bool_vector_to_rows",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_3x1_bool"
+      },
+      "inputs": {
+        "input": { "dtype": "bool", "shape": [3], "data": { "kind": "values", "values": [0, 1, 0] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "bool",
+          "shape": [3, 3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0, 1, 0, 0, 1, 0, 0, 1, 0] }
+        }
+      }
+    },
+    {
+      "name": "ort_bool_col_to_cols",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_1x3_bool",
+        "notes": "The requested shape has a leading dimension of 1; ONNX Expand keeps the input's leading dimension of 3 rather than shrinking it."
+      },
+      "inputs": {
+        "input": { "dtype": "bool", "shape": [3, 1], "data": { "kind": "values", "values": [0, 1, 0] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [1, 3] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "bool",
+          "shape": [3, 3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0, 0, 0, 1, 1, 1, 0, 0, 0] }
+        }
+      }
+    },
+    {
+      "name": "ort_bool_row_to_four_rows",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/expand_test.cc",
+        "test": "ExpandOpTest.Expand_4x1_bool"
+      },
+      "inputs": {
+        "input": { "dtype": "bool", "shape": [1, 4], "data": { "kind": "values", "values": [0, 1, 0, 0] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [4, 1] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "bool",
+          "shape": [4, 4],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0, 1, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0, 0, 1, 0, 0] }
+        }
+      }
+    },
+    {
+      "name": "int8_vec4_path_edge_values",
+      "inputs": {
+        "input": { "dtype": "int8", "shape": [1, 4], "data": { "kind": "values", "values": [-128, -1, 0, 127] } },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 4] } }
+      },
+      "outputs": { "output": { "dtype": "int8", "shape": [3, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "uint8_vec4_path_broadcast_row",
+      "inputs": {
+        "input": {
+          "dtype": "uint8",
+          "shape": [1, 8],
+          "data": { "kind": "values", "values": [0, 64, 128, 192, 255, 1, 127, 254] }
+        },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [4, 8] } }
+      },
+      "outputs": { "output": { "dtype": "uint8", "shape": [4, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "rank7_vec4_last_dim_divisible",
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 1, 3, 1, 2, 1, 4],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.19 }
+        },
+        "shape": { "dtype": "uint32", "shape": [7], "data": { "kind": "values", "values": [2, 3, 3, 2, 2, 2, 4] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 3, 2, 2, 2, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank8_broadcast",
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [1, 2, 1, 3, 1, 1, 2, 1],
+          "data": { "kind": "fillFloat32", "sinStep": 0.07, "cosStep": 0.19 }
+        },
+        "shape": { "dtype": "uint32", "shape": [8], "data": { "kind": "values", "values": [2, 2, 4, 3, 2, 3, 2, 2] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 2, 4, 3, 2, 3, 2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "f16_vec4_row_broadcast",
+      "provenance": {
+        "source": "authored for render coverage",
+        "test": "expand-vec4 `usesF16` arm",
+        "notes": "A float16 output whose innermost dimension is four exercises vec4 broadcasting and half-precision storage. Expand only copies exactly representable values, so the expected output is exact at tolerance zero."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float16",
+          "shape": [1, 8],
+          "data": { "kind": "values", "values": [1.0, -2.0, 3.5, -4.25, 8.0, -16.5, 0.125, 64.0] }
+        },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [3, 8] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [3, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "same_shape_identity_vec4",
+      "provenance": {
+        "source": "authored for render coverage",
+        "test": "expand-vec4 `x_same` arm",
+        "notes": "The requested shape equals the input shape, so expand-vec4 takes the x_same branch and emits the identity offset `return out_index;` instead of including expand-broadcast-offset. No other case has an input shape identical to the output shape while the innermost dim is divisible by 4; every other vec4-eligible case broadcasts at least one axis. All eight values are distinct so a wrong offset in the identity path shows up per position."
+      },
+      "inputs": {
+        "input": {
+          "dtype": "float32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [1.0, -2.0, 3.0, -4.0, 5.0, -6.0, 7.0, -8.0] }
+        },
+        "shape": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [2, 4] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 4], "tolerance": 0 } }
+    }
+  ]
+}