sync 2e7068faf55e

Browse files

Files changed (8) hide show

README.md +68 -0
build/webgpu/bench.json +51 -0
build/webgpu/eyelike-clear-vec4.wgsl.jinja +58 -0
build/webgpu/eyelike-diagonal.wgsl.jinja +46 -0
build/webgpu/eyelike.wgsl.jinja +45 -0
build/webgpu/manifest.json +146 -0
build/webgpu/metadata.json +20 -0
build/webgpu/test.json +316 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.EyeLike
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 22
+## Description
+Generates a 2D identity-like matrix with ones on (or offset from) the main diagonal and zeros everywhere else. The output has the same shape as the 2D input tensor; the output dtype defaults to the input dtype but can be overridden. Attribute `k` shifts the populated diagonal: `k=0` is the main diagonal, `k>0` is upper, `k<0` is lower.
+See the [ONNX `EyeLike` spec](https://onnx.ai/onnx/operators/onnx__EyeLike.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `input` | `input` | `T1` | `2` | — | 2D input tensor whose shape (and optionally type) is copied. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `output` | `output` | `T2` | same as `input` | same as `input` | Output tensor of the same shape as the input, with ones on the selected diagonal and zeros elsewhere. | required |
+## Attributes
+Attributes and default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `k` | `0` | Index of the diagonal to populate with ones: `0` is the main diagonal, positive values select upper diagonals, negative values select lower diagonals. |
+| `dtype` | — | Optional TensorProto DataType enum for the output. When omitted, the output dtype is the same as the input dtype. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T1` | `float32`, `float16`, `uint32`, `int32`, `int16`, `uint8`, `int8`, `bool` |
+| `T2` | `float32`, `float16`, `uint32`, `int32`, `int16`, `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
+- [`eyelike-clear-vec4.wgsl.jinja`](build/webgpu/eyelike-clear-vec4.wgsl.jinja)
+- [`eyelike-diagonal.wgsl.jinja`](build/webgpu/eyelike-diagonal.wgsl.jinja)
+- [`eyelike.wgsl.jinja`](build/webgpu/eyelike.wgsl.jinja)
+## Use with `@huggingface/kernels`
+The loader derives every required output's shape and logical dtype from the manifest contract and this call.
+It then allocates the result tensors automatically.
+The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+Replace each `*Data` placeholder with a typed array containing the corresponding input data.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/ai.onnx.EyeLike", { version: 1 });
+const { output } = await kernel({ input: { data: inputData, shape: [2, 2] } });
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,51 @@

+{
+  "op": "ai.onnx.EyeLike",
+  "cases": [
+    {
+      "name": "eyelike-f32-2048-square",
+      "inputs": { "input": { "dtype": "float32", "shape": [2048, 2048] } },
+      "outputs": { "output": { "dtype": "float32", "shape": [2048, 2048] } }
+    },
+    {
+      "name": "eyelike-f32-4097-square-tail-healthy",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "count": 16785409 },
+      "inputs": { "input": { "dtype": "float32", "shape": [4097, 4097] } },
+      "outputs": { "output": { "dtype": "float32", "shape": [4097, 4097] } },
+      "bench": { "primary": true, "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype)" }] }
+    },
+    {
+      "name": "eyelike-f16-4096-square-narrow-store",
+      "preset": "smoke",
+      "vars": { "dtype": "float16", "count": 16777216 },
+      "inputs": { "input": { "dtype": "float16", "shape": [4096, 4096] } },
+      "outputs": { "output": { "dtype": "float16", "shape": [4096, 4096] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype)" }] }
+    },
+    {
+      "name": "eyelike-int8-4096-square-word-store-pathology",
+      "preset": "smoke",
+      "vars": { "dtype": "int8", "count": 16777216 },
+      "inputs": { "input": { "dtype": "int8", "shape": [4096, 4096] } },
+      "outputs": { "output": { "dtype": "int8", "shape": [4096, 4096] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * 4" }] }
+    },
+    {
+      "name": "eyelike-f32-134m-2dfold-dispatch-cliff",
+      "preset": "stress",
+      "provenance": { "notes": "Stress-only capacity case: the input and output together occupy 1 GiB of GPU storage." },
+      "vars": { "dtype": "float32", "count": 134217728 },
+      "inputs": { "input": { "dtype": "float32", "shape": [16384, 8192] } },
+      "outputs": { "output": { "dtype": "float32", "shape": [16384, 8192] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype)" }] }
+    },
+    {
+      "name": "eyelike-f32-8192-square-long-context-mask",
+      "preset": "stress",
+      "vars": { "dtype": "float32", "count": 67108864 },
+      "inputs": { "input": { "dtype": "float32", "shape": [8192, 8192] } },
+      "outputs": { "output": { "dtype": "float32", "shape": [8192, 8192] } },
+      "bench": { "primary": true, "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype)" }] }
+    }
+  ]
+}

build/webgpu/eyelike-clear-vec4.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,58 @@

+{% 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 COUNT4: u32 = {{ source.count4 }}u;
+{% if tailSafe %}
+const COUNT: u32 = {{ source.count }}u;
+{% endif %}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d("q", "", note="") }}
+  if (q < COUNT4) {
+{% if tailSafe %}
+    let base = q * 4u;
+    output[base] = {{ outScalar }}(0);
+    if (base + 1u < COUNT) { output[base + 1u] = {{ outScalar }}(0); }
+    if (base + 2u < COUNT) { output[base + 2u] = {{ outScalar }}(0); }
+    if (base + 3u < COUNT) { output[base + 3u] = {{ outScalar }}(0); }
+{% else %}
+    output[q] = {{ outVector }}({{ outScalar }}(0));
+{% endif %}
+  }
+}

build/webgpu/eyelike-diagonal.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,46 @@

+{% 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 }}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d("row", "params.rows", guardInline=true, note="") }}
+  let col = i32(row) + params.k;
+  if (col >= 0 && col < i32(params.cols)) {
+    output[row * params.cols + u32(col)] = {{ outScalar }}(1);
+  }
+}

build/webgpu/eyelike.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,45 @@

+{% 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 }}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d() }}
+  let row = i32(i / params.cols);
+  let col = i32(i % params.cols);
+  output[i] = select({{ outScalar }}(0), {{ outScalar }}(1), col - row == params.k);
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,146 @@

+{
+  "domain": "ai.onnx",
+  "name": "EyeLike",
+  "sinceVersion": 22,
+  "description": "Generates a 2D identity-like matrix with ones on (or offset from) the main diagonal and zeros everywhere else. The output has the same shape as the 2D input tensor; the output dtype defaults to the input dtype but can be overridden. Attribute `k` shifts the populated diagonal: `k=0` is the main diagonal, `k>0` is upper, `k<0` is lower.",
+  "inputs": [
+    {
+      "role": "input",
+      "dtype": "T1",
+      "rank": 2,
+      "description": "2D input tensor whose shape (and optionally type) is copied."
+    }
+  ],
+  "outputs": [
+    {
+      "role": "output",
+      "dtype": "T2",
+      "rank": "ranks.input",
+      "description": "Output tensor of the same shape as the input, with ones on the selected diagonal and zeros elsewhere.",
+      "shape": "shapes.input"
+    }
+  ],
+  "attributes": { "k": 0 },
+  "attributeDescriptions": {
+    "dtype": "Optional TensorProto DataType enum for the output. When omitted, the output dtype is the same as the input dtype.",
+    "k": "Index of the diagonal to populate with ones: `0` is the main diagonal, positive values select upper diagonals, negative values select lower diagonals."
+  },
+  "typeConstraints": {
+    "T1": ["float32", "float16", "uint32", "int32", "int16", "uint8", "int8", "bool"],
+    "T2": ["float32", "float16", "uint32", "int32", "int16", "uint8", "int8", "bool"]
+  },
+  "args": {
+    "input": { "kind": "tensor", "semantic": "input", "role": "input" },
+    "output": { "kind": "tensor", "semantic": "output", "role": "output" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "derive": {
+    "wave32Adapter": "has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize == 32 and device.adapterInfo.subgroupMaxSize == 32",
+    "reportedNonWave32Adapter": "not wave32Adapter and (has(device.adapterInfo, \"subgroupMinSize\") or has(device.adapterInfo, \"subgroupMaxSize\"))",
+    "dtypeContract": "(has(attrs, \"dtype\") and attrs.dtype == onnxDtypeCode(logicalDtypes.T2)) or (not has(attrs, \"dtype\") and logicalDtypes.T2 == logicalDtypes.T1)",
+    "packedClearPreferred": "numel(shapes.output) % 4 == 0 or not reportedNonWave32Adapter"
+  },
+  "constants": { "usesF16": "tensorDtypes.output == \"float16\"", "outScalar": "dtypes.T2" },
+  "variants": [
+    {
+      "id": "rank2_vec4",
+      "priority": 20,
+      "when": ["dtypeContract", "ranks.input == 2", "ranks.output == 2", "dim(shapes.output, 0) == dim(shapes.input, 0)", "dim(shapes.output, 1) == dim(shapes.input, 1)", "dim(shapes.output, 1) > 0", "packedClearPreferred", "f16Ok(tensorDtypes.output)"],
+      "constants": {
+        "tailSafe": "numel(shapes.output) % 4 != 0",
+        "outVector": "\"vec4<\" ~ dtypes.T2 ~ \">\"",
+        "clearElement": "dtypes.T2 if numel(shapes.output) % 4 != 0 else (\"vec4<\" ~ dtypes.T2 ~ \">\")"
+      },
+      "passes": [
+        {
+          "id": "clear",
+          "name": "EyeLike.ClearVec4",
+          "source": {
+            "shader": "eyelike-clear-vec4.wgsl.jinja",
+            "inputs": { "count4": "ceilDiv(numel(shapes.output), 4)", "count": "numel(shapes.output)" }
+          },
+          "bindings": [
+            {
+              "name": "output",
+              "arg": "output",
+              "semantic": "output",
+              "buffer": { "type": "storage" },
+              "elementType": "$clearElement"
+            }
+          ],
+          "dispatch": { "threads": "ceilDiv(numel(shapes.output), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        },
+        {
+          "id": "diagonal",
+          "name": "EyeLike.Diagonal",
+          "shader": "eyelike-diagonal.wgsl.jinja",
+          "bindings": [
+            {
+              "name": "output",
+              "arg": "output",
+              "semantic": "output",
+              "buffer": { "type": "storage" },
+              "elementType": "$outScalar"
+            },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [
+                  { "name": "rows", "type": "u32", "value": "dim(shapes.output, 0)" },
+                  { "name": "cols", "type": "u32", "value": "dim(shapes.output, 1)" },
+                  {
+                    "name": "k",
+                    "type": "i32",
+                    "value": "max(min(attrs.k, dim(shapes.output, 1)), 0 - dim(shapes.output, 0))"
+                  }
+                ]
+              }
+            }
+          ],
+          "dispatch": { "threads": "dim(shapes.output, 0)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "rank2",
+      "when": ["dtypeContract", "ranks.input == 2", "ranks.output == 2", "dim(shapes.output, 0) == dim(shapes.input, 0)", "dim(shapes.output, 1) == dim(shapes.input, 1)", "f16Ok(tensorDtypes.output)"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "EyeLike",
+          "shader": "eyelike.wgsl.jinja",
+          "bindings": [
+            {
+              "name": "output",
+              "arg": "output",
+              "semantic": "output",
+              "buffer": { "type": "storage" },
+              "elementType": "$outScalar"
+            },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [
+                  { "name": "count", "type": "u32", "value": "numel(shapes.output)" },
+                  { "name": "cols", "type": "u32", "value": "dim(shapes.output, 1)" },
+                  {
+                    "name": "k",
+                    "type": "i32",
+                    "value": "max(min(attrs.k, dim(shapes.output, 1)), 0 - dim(shapes.output, 0))"
+                  }
+                ]
+              }
+            }
+          ],
+          "dispatch": { "threads": "numel(shapes.output)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,20 @@

+{
+  "name": "ai.onnx.EyeLike",
+  "id": "_ai_onnx_eyelike_webgpu_abf2012",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "AT1FVETPSt5mKqAQTjxmQ6hYC4cGMKFP6rnonySkJfE=",
+      "eyelike-clear-vec4.wgsl.jinja": "J0mLIcOnYIZzMcidLRWN1BS9xb1fqIeJ34dHmc+dqII=",
+      "eyelike-diagonal.wgsl.jinja": "mDunKMVacGusMV5mH1aLeddRfAffsCA0wwymCcCzVtg=",
+      "eyelike.wgsl.jinja": "G7fUoLpUu4wd3zcg6qsQOOYRVUKCduoTf1Rs3Tg2xUA=",
+      "manifest.json": "JtsNde7hNNptyv5Z2eQWDU/QAtUzbt8RmnOjd0a6J7A=",
+      "test.json": "+NBKAgSG911NlnbXaGR9KwZCbiNT3XHpsYT8t1pHWuQ="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.EyeLike" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,316 @@

+{
+  "op": "ai.onnx.EyeLike",
+  "cases": [
+    {
+      "name": "vec4_upper_diagonal_rectangular_f32",
+      "attrs": { "k": 1 },
+      "inputs": { "input": { "dtype": "float32", "shape": [4, 8], "data": { "kind": "constant", "value": 7.0 } } },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "dispatch_cliff_f32_tall",
+      "attrs": { "k": 0 },
+      "inputs": {
+        "input": { "dtype": "float32", "shape": [16777217, 1], "data": { "kind": "constant", "value": 0.0 } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [16777217, 1], "tolerance": 0 } }
+    },
+    {
+      "name": "default_f32_rectangular",
+      "inputs": { "input": { "dtype": "float32", "shape": [3, 2], "data": { "kind": "constant", "value": 0.0 } } },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 2] } },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/eyelike_op_test.cc",
+        "test": "EyeLikeOpTest.EyeLikeDefault"
+      }
+    },
+    {
+      "name": "zero_rows_f32_empty_output",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/eyelike_op_test.cc",
+        "test": "EyeLikeOpTest.EyeLikeDefault",
+        "notes": "Zero-row rank-2 inputs produce an empty output; no diagonal elements are materialized."
+      },
+      "inputs": { "input": { "dtype": "float32", "shape": [0, 3], "data": { "kind": "values", "values": [] } } },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [0, 3], "tolerance": 0, "data": { "kind": "values", "values": [] } }
+      }
+    },
+    {
+      "name": "upper_diagonal_f32",
+      "attrs": { "k": 2 },
+      "inputs": { "input": { "dtype": "float32", "shape": [3, 4], "data": { "kind": "constant", "value": 0.0 } } },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 4] } },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/eyelike_op_test.cc",
+        "test": "EyeLikeOpTest.EyeLike_UpperDiagonal"
+      }
+    },
+    {
+      "name": "lower_diagonal_f32",
+      "attrs": { "k": -1 },
+      "inputs": { "input": { "dtype": "float32", "shape": [3, 2], "data": { "kind": "constant", "value": 0.0 } } },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 2] } },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/eyelike_op_test.cc",
+        "test": "EyeLikeOpTest.EyeLike_LowerDiagonal"
+      }
+    },
+    {
+      "name": "out_of_range_int32",
+      "attrs": { "k": 3 },
+      "inputs": { "input": { "dtype": "int32", "shape": [3, 2], "data": { "kind": "constant", "value": 0 } } },
+      "outputs": { "output": { "dtype": "int32", "shape": [3, 2] } },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/eyelike_op_test.cc",
+        "test": "EyeLikeOpTest.EyeLike_K_EgdeCase_1",
+        "notes": "ORT uses int64 dtype; this framework covers the same out-of-range positive diagonal with int32 output storage."
+      }
+    },
+    {
+      "name": "ort_out_of_range_negative_int32",
+      "attrs": { "k": -3 },
+      "inputs": { "input": { "dtype": "int32", "shape": [3, 2], "data": { "kind": "constant", "value": 0 } } },
+      "outputs": { "output": { "dtype": "int32", "shape": [3, 2] } },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/eyelike_op_test.cc",
+        "test": "EyeLikeOpTest.EyeLike_K_EgdeCase_2",
+        "notes": "ORT uses int64 dtype; this framework covers the same out-of-range negative diagonal with int32 output storage."
+      }
+    },
+    {
+      "name": "large_positive_k_does_not_wrap_to_main_diagonal",
+      "attrs": { "k": 4294967296 },
+      "inputs": { "input": { "dtype": "float32", "shape": [3, 3], "data": { "kind": "constant", "value": 0.0 } } },
+      "outputs": {
+        "output": { "dtype": "float32", "shape": [3, 3], "tolerance": 0, "data": { "kind": "constant", "value": 0.0 } }
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/eyelike_op_test.cc",
+        "test": "EyeLikeOpTest.EyeLike_K_EgdeCase_1",
+        "notes": "ONNX k is an int64 attribute. A huge positive offset is valid and should be out of range, not wrap through a 32-bit uniform to the main diagonal."
+      }
+    },
+    {
+      "name": "large_negative_k_does_not_wrap_to_main_diagonal",
+      "attrs": { "k": -4294967296 },
+      "inputs": { "input": { "dtype": "int32", "shape": [3, 3], "data": { "kind": "constant", "value": 0 } } },
+      "outputs": {
+        "output": { "dtype": "int32", "shape": [3, 3], "tolerance": 0, "data": { "kind": "constant", "value": 0 } }
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/eyelike_op_test.cc",
+        "test": "EyeLikeOpTest.EyeLike_K_EgdeCase_2",
+        "notes": "Negative companion for int64 k values outside the int32 range; the diagonal remains out of range."
+      }
+    },
+    {
+      "name": "main_diagonal_f16",
+      "attrs": { "dtype": 10 },
+      "inputs": { "input": { "dtype": "float16", "shape": [2, 3], "data": { "kind": "constant", "value": 0.0 } } },
+      "outputs": { "output": { "dtype": "float16", "shape": [2, 3] } }
+    },
+    {
+      "name": "uint8_output_main_diagonal",
+      "attrs": { "dtype": 2 },
+      "inputs": { "input": { "dtype": "uint8", "shape": [2, 4], "data": { "kind": "constant", "value": 99 } } },
+      "outputs": { "output": { "dtype": "uint8", "shape": [2, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "int8_output_lower_diagonal",
+      "attrs": { "k": -1, "dtype": 3 },
+      "inputs": { "input": { "dtype": "int8", "shape": [4, 3], "data": { "kind": "constant", "value": -7 } } },
+      "outputs": { "output": { "dtype": "int8", "shape": [4, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "uint32_extreme_input_main_diagonal",
+      "provenance": {
+        "source": "onnx/onnx/docs/Operators.md#EyeLike-22",
+        "notes": "Covers standard uint32 input and output membership; extreme input payloads are intentionally ignored because EyeLike consumes only the input shape."
+      },
+      "attrs": { "dtype": 12 },
+      "inputs": {
+        "input": {
+          "dtype": "uint32",
+          "shape": [2, 4],
+          "data": { "kind": "values", "values": [0, 4294967295, 2147483648, 16777217, 1, 42, 4000000000, 255] }
+        }
+      },
+      "outputs": { "output": { "dtype": "uint32", "shape": [2, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "int16_extreme_input_lower_diagonal",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/eyelike_op_test.cc",
+        "test": "EyeLikeOpTest.EyeLike_LowerDiagonal",
+        "notes": "Extends ORT's lower-diagonal case to ONNX-standard int16 and proves that only the input shape, not its extreme values, affects the result."
+      },
+      "attrs": { "k": -1, "dtype": 5 },
+      "inputs": {
+        "input": { "dtype": "int16", "shape": [3, 5], "data": { "kind": "cycle", "values": [-32768, 32767, -1, 0, 1] } }
+      },
+      "outputs": { "output": { "dtype": "int16", "shape": [3, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "bool_output_upper_diagonal",
+      "provenance": {
+        "source": "onnx/onnx/docs/Operators.md#EyeLike-22",
+        "notes": "Covers the ONNX-standard bool input/output route with a non-default upper diagonal."
+      },
+      "attrs": { "k": 1, "dtype": 9 },
+      "inputs": { "input": { "dtype": "bool", "shape": [3, 4], "data": { "kind": "cycle", "values": [0, 1, 1, 0] } } },
+      "outputs": { "output": { "dtype": "bool", "shape": [3, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "upper_diagonal_k1_narrow",
+      "attrs": { "k": 1, "dtype": 2 },
+      "inputs": { "input": { "dtype": "float32", "shape": [3, 2], "data": { "kind": "constant", "value": 0.0 } } },
+      "outputs": { "output": { "dtype": "uint8", "shape": [3, 2], "tolerance": 0 } },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/eyelike_op_test.cc",
+        "test": "EyeLikeOpTest.EyeLike_UpperrDiagonal2",
+        "notes": "The same narrow upper-diagonal shape also covers the independent uint8 output-dtype route."
+      }
+    },
+    {
+      "name": "lower_diagonal_k_minus2_wide",
+      "attrs": { "k": -2, "dtype": 3 },
+      "inputs": { "input": { "dtype": "float32", "shape": [3, 4], "data": { "kind": "constant", "value": 0.0 } } },
+      "outputs": { "output": { "dtype": "int8", "shape": [3, 4], "tolerance": 0 } },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/eyelike_op_test.cc",
+        "test": "EyeLikeOpTest.EyeLike_LowerDiagonal2",
+        "notes": "The same wide lower-diagonal shape also covers the independent int8 output-dtype route."
+      }
+    },
+    {
+      "name": "ort_different_output_dtype_int32",
+      "attrs": { "dtype": 6 },
+      "inputs": { "input": { "dtype": "float32", "shape": [3, 3], "data": { "kind": "constant", "value": 0.0 } } },
+      "outputs": { "output": { "dtype": "int32", "shape": [3, 3], "tolerance": 0 } },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/eyelike_op_test.cc",
+        "test": "EyeLikeOpTest.EyeLike_DifferentDtype",
+        "notes": "ORT requests int64 via dtype attr; this framework expresses the same output-dtype behavior with supported int32 storage."
+      }
+    },
+    {
+      "name": "int32_output_nonzero_input_ignored",
+      "inputs": { "input": { "dtype": "int32", "shape": [2, 2], "data": { "kind": "constant", "value": 99 } } },
+      "outputs": { "output": { "dtype": "int32", "shape": [2, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_eyelike_populate_off_main_diagonal",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_eyelike_populate_off_main_diagonal"
+      },
+      "attrs": { "k": 1, "dtype": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "int32",
+          "shape": [4, 5],
+          "data": {
+            "kind": "values",
+            "values": [44, 47, 64, 67, 67, 9, 83, 21, 36, 87, 70, 88, 88, 12, 58, 65, 39, 87, 46, 88]
+          }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_eyelike_with_dtype",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_eyelike_with_dtype",
+        "notes": "Official output dtype is DOUBLE; adapted to FLOAT because this manifest does not support float64 outputs."
+      },
+      "attrs": { "dtype": 1 },
+      "inputs": {
+        "input": {
+          "dtype": "int32",
+          "shape": [3, 4],
+          "data": { "kind": "values", "values": [44, 47, 64, 67, 67, 9, 83, 21, 36, 87, 70, 88] }
+        }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_eyelike_without_dtype",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_eyelike_without_dtype" },
+      "attrs": {},
+      "inputs": {
+        "input": {
+          "dtype": "int32",
+          "shape": [4, 4],
+          "data": { "kind": "values", "values": [44, 47, 64, 67, 67, 9, 83, 21, 36, 87, 70, 88, 88, 12, 58, 65] }
+        }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "empty_input_zero_dim",
+      "inputs": { "input": { "dtype": "float32", "shape": [0, 2], "data": { "kind": "values", "values": [] } } },
+      "outputs": { "output": { "dtype": "float32", "shape": [0, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "f32_input_f16_output_64x64_upper_diagonal_k1",
+      "provenance": {
+        "notes": "Separates the float16 output route from the independently typed float32 shape-source input."
+      },
+      "attrs": { "k": 1, "dtype": 10 },
+      "inputs": { "input": { "dtype": "float32", "shape": [64, 64], "data": { "kind": "constant", "value": 0.0 } } },
+      "outputs": { "output": { "dtype": "float16", "shape": [64, 64], "tolerance": 0 } }
+    },
+    {
+      "name": "int32_nonsquare_tall_lower_diagonal_k_minus2",
+      "attrs": { "k": -2 },
+      "inputs": { "input": { "dtype": "int32", "shape": [5, 3], "data": { "kind": "constant", "value": 0 } } },
+      "outputs": { "output": { "dtype": "int32", "shape": [5, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "k_at_negative_rows_boundary_all_zeros",
+      "provenance": {
+        "notes": "A float32 shape-source input with bool output covers the independent bool dtype route on an all-off diagonal."
+      },
+      "attrs": { "k": -3, "dtype": 9 },
+      "inputs": { "input": { "dtype": "float32", "shape": [3, 3], "data": { "kind": "constant", "value": 0.0 } } },
+      "outputs": {
+        "output": { "dtype": "bool", "shape": [3, 3], "tolerance": 0, "data": { "kind": "constant", "value": 0 } }
+      }
+    },
+    {
+      "name": "k_zero_1x1_single_element",
+      "provenance": {
+        "notes": "A float32 shape-source input with int16 output covers the independent int16 dtype route."
+      },
+      "attrs": { "k": 0, "dtype": 5 },
+      "inputs": { "input": { "dtype": "float32", "shape": [1, 1], "data": { "kind": "constant", "value": 0.0 } } },
+      "outputs": {
+        "output": { "dtype": "int16", "shape": [1, 1], "tolerance": 0, "data": { "kind": "values", "values": [1] } }
+      }
+    },
+    {
+      "name": "k_at_positive_cols_boundary_all_zeros",
+      "provenance": {
+        "notes": "A float32 shape-source input with uint32 output covers the independent uint32 dtype route on an all-off diagonal."
+      },
+      "attrs": { "k": 3, "dtype": 12 },
+      "inputs": { "input": { "dtype": "float32", "shape": [4, 3], "data": { "kind": "constant", "value": 0.0 } } },
+      "outputs": {
+        "output": { "dtype": "uint32", "shape": [4, 3], "tolerance": 0, "data": { "kind": "constant", "value": 0 } }
+      }
+    },
+    {
+      "name": "k_minus_rows_minus1_single_corner",
+      "attrs": { "k": -3 },
+      "inputs": { "input": { "dtype": "float32", "shape": [4, 3], "data": { "kind": "constant", "value": 0.0 } } },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [4, 3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 0.0, 1.0, 0.0, 0.0] }
+        }
+      }
+    }
+  ]
+}