sync 2e7068faf55e

Browse files

Files changed (8) hide show

README.md +82 -0
build/webgpu/bench.json +90 -0
build/webgpu/manifest.json +278 -0
build/webgpu/metadata.json +20 -0
build/webgpu/one-hot-fill.wgsl.jinja +50 -0
build/webgpu/one-hot-last-axis-vec4.wgsl.jinja +117 -0
build/webgpu/one-hot-scatter.wgsl.jinja +69 -0
build/webgpu/test.json +790 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.OneHot
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 11
+## Description
+Produces a one-hot tensor from an `indices` input: positions matching each index are filled with `on_value` and all other positions with `off_value`, where both are taken from the two-element `values` tensor `[off_value, on_value]`. The output rank is one greater than `indices`, with the new dimension of size `depth` inserted at the position given by `axis`; indices outside `[-depth, depth-1]` yield all-`off_value` rows.
+See the [ONNX `OneHot` spec](https://onnx.ai/onnx/operators/onnx__OneHot.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `indices` | `indices` | `I` | — | — | Integer or float index tensor; values outside `[-depth, depth-1]` produce all-`off_value` output rows. | required |
+| `depth` | `depth` | `D` | — | — | Scalar (or length-1 rank-1) tensor specifying the number of classes and the size of the one-hot dimension. | required |
+| `values` | `values` | `T` | `1` | — | Rank-1 tensor of exactly two elements `[off_value, on_value]` giving the values written to inactive and active positions respectively. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `output` | `output` | `T` | derived | — | One-hot tensor with rank equal to `rank(indices) + 1`, same element type as `values`. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `axis` | `-1` | Axis along which the one-hot dimension is inserted; default `-1` appends it as the last dimension. Negative values count from the back; accepted range is `[-r-1, r]` where `r = rank(indices)`. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `I` | `float32`, `float16`, `int32`, `int16`, `int8`, `uint32`, `uint8` |
+| `D` | `float32`, `float16`, `int32`, `int16`, `int8`, `uint32`, `uint8` |
+| `T` | `float32`, `float16`, `int32`, `int16`, `int8`, `uint32`, `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
+- [`one-hot-fill.wgsl.jinja`](build/webgpu/one-hot-fill.wgsl.jinja)
+- [`one-hot-last-axis-vec4.wgsl.jinja`](build/webgpu/one-hot-last-axis-vec4.wgsl.jinja)
+- [`one-hot-scatter.wgsl.jinja`](build/webgpu/one-hot-scatter.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.OneHot", { version: 1 });
+// Explicit destinations request optional results or supply metadata that cannot be inferred.
+const { output } = await kernel({
+  indices: { data: indicesData, shape: [1] },
+  depth: { data: depthData, shape: [] },
+  values: { data: valuesData, shape: [2] },
+}, {
+  outputs: { output: { shape: [1, 2], dtype: "float32" } },
+});
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,90 @@

+{
+  "op": "ai.onnx.OneHot",
+  "cases": [
+    {
+      "name": "tokens_vocab",
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [128] },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [4096] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [128, 4096] } }
+    },
+    {
+      "name": "tokens_vocab_2047x8192",
+      "preset": "smoke",
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [2047], "dist": "uniform", "seed": 901, "scale": 8000, "offset": 0 },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [8192] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2047, 8192] } },
+      "bench": { "primary": true, "metrics": [{ "type": "bandwidth", "value": "numel(shapes.output) * 4" }] }
+    },
+    {
+      "name": "unaligned_depth_8191_generic_scalar",
+      "preset": "smoke",
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [2047], "dist": "uniform", "seed": 902, "scale": 8000, "offset": 0 },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [8191] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2047, 8191] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "numel(shapes.output) * 4" }] }
+    },
+    {
+      "name": "axis0_generic_8192x2047",
+      "preset": "smoke",
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [2047], "dist": "uniform", "seed": 903, "scale": 8000, "offset": 0 },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [8192] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [8192, 2047] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "numel(shapes.output) * 4" }] }
+    },
+    {
+      "name": "axis1_middle_generic_2047x1024x8",
+      "preset": "smoke",
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [2047, 8], "dist": "uniform", "seed": 904, "scale": 1000, "offset": 0 },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1024] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2047, 1024, 8] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "numel(shapes.output) * 4" }] }
+    },
+    {
+      "name": "tokens_vocab_f16_2047x8192_vec4",
+      "preset": "smoke",
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [2047], "dist": "uniform", "seed": 905, "scale": 8000, "offset": 0 },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [8192] } },
+        "values": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [2047, 8192] } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "numel(shapes.output) * 2" }] }
+    },
+    {
+      "name": "generic_middle_axis_rank5_3p9m",
+      "preset": "stress",
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "indices": {
+          "dtype": "uint32",
+          "shape": [256, 8, 8, 8],
+          "dist": "uniform",
+          "seed": 911,
+          "scale": 29,
+          "offset": 0
+        },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [30] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [256, 30, 8, 8, 8], "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "numel(shapes.output) * 4" }] }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,278 @@

+{
+  "domain": "ai.onnx",
+  "name": "OneHot",
+  "sinceVersion": 11,
+  "description": "Produces a one-hot tensor from an `indices` input: positions matching each index are filled with `on_value` and all other positions with `off_value`, where both are taken from the two-element `values` tensor `[off_value, on_value]`. The output rank is one greater than `indices`, with the new dimension of size `depth` inserted at the position given by `axis`; indices outside `[-depth, depth-1]` yield all-`off_value` rows.",
+  "inputs": [
+    {
+      "role": "indices",
+      "dtype": "I",
+      "description": "Integer or float index tensor; values outside `[-depth, depth-1]` produce all-`off_value` output rows."
+    },
+    {
+      "role": "depth",
+      "dtype": "D",
+      "description": "Scalar (or length-1 rank-1) tensor specifying the number of classes and the size of the one-hot dimension."
+    },
+    {
+      "role": "values",
+      "dtype": "T",
+      "description": "Rank-1 tensor of exactly two elements `[off_value, on_value]` giving the values written to inactive and active positions respectively.",
+      "rank": 1
+    }
+  ],
+  "outputs": [
+    {
+      "role": "output",
+      "dtype": "T",
+      "description": "One-hot tensor with rank equal to `rank(indices) + 1`, same element type as `values`.",
+      "rank": "ranks.indices + 1"
+    }
+  ],
+  "attributes": { "axis": -1 },
+  "attributeDescriptions": {
+    "axis": "Axis along which the one-hot dimension is inserted; default `-1` appends it as the last dimension. Negative values count from the back; accepted range is `[-r-1, r]` where `r = rank(indices)`."
+  },
+  "typeConstraints": {
+    "I": ["float32", "float16", "int32", "int16", "int8", "uint32", "uint8"],
+    "D": ["float32", "float16", "int32", "int16", "int8", "uint32", "uint8"],
+    "T": ["float32", "float16", "int32", "int16", "int8", "uint32", "uint8", "bool"]
+  },
+  "args": {
+    "indices": { "kind": "tensor", "semantic": "indices", "role": "indices" },
+    "depth": { "kind": "tensor", "semantic": "depth", "role": "depth" },
+    "values": { "kind": "tensor", "semantic": "values", "role": "values" },
+    "output": { "kind": "tensor", "semantic": "output", "role": "output" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "derive": {
+    "oneHotContractOk": "(dtypes.I == \"u32\" or dtypes.I == \"i32\" or dtypes.I == \"f32\" or dtypes.I == \"f16\") and ranks.indices >= 1 and ranks.output == ranks.indices + 1 and (ranks.depth == 0 or (ranks.depth == 1 and dim(shapes.depth, 0) == 1)) and ranks.values == 1 and dim(shapes.values, 0) == 2 and numel(shapes.output) >= numel(shapes.indices) and f16Ok(dtypes.T) and f16Ok(dtypes.I)",
+    "oneHotLastAxisCovered": "oneHotContractOk and (attrs.axis == ranks.output - 1 or attrs.axis == -1) and dim(shapes.output, -1) % 4 == 0",
+    "oneHotLastAxisVecCount": "dim(shapes.output, -1) / 4",
+    "oneHotRowWorkgroup": "min(tunables.WORKGROUP_SIZE, device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX, max(1, pow2ceil(oneHotLastAxisVecCount)))",
+    "oneHotCooperativeMinVecs": "device.adapterInfo.subgroupMinSize if has(device.adapterInfo, \"subgroupMinSize\") else 32",
+    "oneHotRowCooperative": "oneHotLastAxisCovered and oneHotLastAxisVecCount >= oneHotCooperativeMinVecs"
+  },
+  "constants": {
+    "scalar": "dtypes.T",
+    "indexScalar": "dtypes.I",
+    "usesI32": "dtypes.I == \"i32\"",
+    "floatIndices": "dtypes.I == \"f32\" or dtypes.I == \"f16\"",
+    "usesF16": "dtypes.T == \"f16\" or dtypes.I == \"f16\"",
+    "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\""
+  },
+  "variants": [
+    {
+      "id": "last_axis_row_vec4",
+      "priority": 20,
+      "when": ["oneHotRowCooperative"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "OneHot.last_axis_row_vec4",
+          "source": {
+            "shader": "one-hot-last-axis-vec4.wgsl.jinja",
+            "inputs": {
+              "depth": "dim(shapes.output, -1)",
+              "rows": "numel(shapes.indices)",
+              "wg": "oneHotRowWorkgroup",
+              "cooperative": true
+            }
+          },
+          "bindings": [
+            {
+              "name": "indices",
+              "arg": "indices",
+              "semantic": "indices",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$indexScalar"
+            },
+            {
+              "name": "values",
+              "arg": "values",
+              "semantic": "values",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar",
+              "length": 2
+            },
+            {
+              "name": "output",
+              "arg": "output",
+              "semantic": "output",
+              "buffer": { "type": "storage" },
+              "elementType": "$vectorScalar"
+            }
+          ],
+          "dispatch": { "workgroups": "numel(shapes.indices)" }
+        }
+      ]
+    },
+    {
+      "id": "last_axis_vec4",
+      "priority": 10,
+      "when": ["oneHotLastAxisCovered"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "OneHot.last_axis_vec4",
+          "source": {
+            "shader": "one-hot-last-axis-vec4.wgsl.jinja",
+            "inputs": { "depth": "dim(shapes.output, -1)", "cooperative": false }
+          },
+          "bindings": [
+            {
+              "name": "indices",
+              "arg": "indices",
+              "semantic": "indices",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$indexScalar"
+            },
+            {
+              "name": "values",
+              "arg": "values",
+              "semantic": "values",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar",
+              "length": 2
+            },
+            {
+              "name": "output",
+              "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": "generic_axis",
+      "when": ["oneHotContractOk", "attrs.axis + ranks.output >= 0", "attrs.axis < ranks.output"],
+      "derive": {
+        "depthAxis": "attrs.axis if attrs.axis >= 0 else attrs.axis + ranks.output",
+        "depthInner": "max(1, inner(shapes.output, depthAxis))",
+        "outputVecCount": "floor(numel(shapes.output) / 4)",
+        "outputTailCount": "numel(shapes.output) % 4"
+      },
+      "passes": [
+        {
+          "id": "fill_vec4",
+          "name": "OneHot.FillVec4",
+          "source": { "shader": "one-hot-fill.wgsl.jinja", "inputs": { "vectorized": true } },
+          "bindings": [
+            {
+              "name": "values",
+              "arg": "values",
+              "semantic": "values",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar",
+              "length": 2
+            },
+            {
+              "name": "output",
+              "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": "outputVecCount" }] }
+            }
+          ],
+          "dispatch": { "threads": "outputVecCount", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        },
+        {
+          "id": "fill_tail",
+          "name": "OneHot.FillTail",
+          "source": { "shader": "one-hot-fill.wgsl.jinja", "inputs": { "vectorized": false } },
+          "bindings": [
+            {
+              "name": "values",
+              "arg": "values",
+              "semantic": "values",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar",
+              "length": 2
+            },
+            {
+              "name": "output",
+              "arg": "output",
+              "semantic": "output",
+              "buffer": { "type": "storage" },
+              "elementType": "$scalar"
+            },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [
+                  { "name": "count", "type": "u32", "value": "outputTailCount" },
+                  { "name": "offset", "type": "u32", "value": "outputVecCount * 4" }
+                ]
+              }
+            }
+          ],
+          "dispatch": { "threads": "outputTailCount", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        },
+        {
+          "id": "scatter",
+          "name": "OneHot.Scatter",
+          "source": {
+            "shader": "one-hot-scatter.wgsl.jinja",
+            "inputs": { "depth": "dim(shapes.output, depthAxis)", "inner": "depthInner" }
+          },
+          "bindings": [
+            {
+              "name": "indices",
+              "arg": "indices",
+              "semantic": "indices",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$indexScalar"
+            },
+            {
+              "name": "values",
+              "arg": "values",
+              "semantic": "values",
+              "buffer": { "type": "read-only-storage" },
+              "elementType": "$scalar",
+              "length": 2
+            },
+            {
+              "name": "output",
+              "arg": "output",
+              "semantic": "output",
+              "buffer": { "type": "storage" },
+              "elementType": "$scalar"
+            },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": {
+                "name": "Params",
+                "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.indices)" }]
+              }
+            }
+          ],
+          "dispatch": { "threads": "numel(shapes.indices)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

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

+{
+  "name": "ai.onnx.OneHot",
+  "id": "_ai_onnx_onehot_webgpu_21894c9",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "8JSemsHPcNMzWPXe0x90can66FWOW5qWvz0P7W7I7Pg=",
+      "manifest.json": "BfEMlX6Fvw/2pFVCJ1L9FoZ2lBpbHjLoAr2PbX2GBZc=",
+      "one-hot-fill.wgsl.jinja": "+1+3T4yVF91Ta11a5kEZHAC8D3C/uGzVKZm/CSK4Htc=",
+      "one-hot-last-axis-vec4.wgsl.jinja": "nQ+vRZPx/rWi3Q+vuIhccDJzW6BATwWLFiaDf+HNwXo=",
+      "one-hot-scatter.wgsl.jinja": "xXJjRoPmpglSeIQvxCv+9mmygMaGMw1LTKy82ViMVJ8=",
+      "test.json": "UrZtlg9pSruj4lfpgeNOargkGctAr+ZUIM0fkeBJNsg="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.OneHot" }
+}

build/webgpu/one-hot-fill.wgsl.jinja ADDED Viewed

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

+{% 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 }}
+// Dense off-value fill for arbitrary-axis OneHot. The vector form writes four
+// elements per invocation and the scalar form handles the tail. A sparse pass
+// writes the on-values afterward.
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d(note="") }}
+{% if source.vectorized %}
+  output[i] = {{ vectorScalar }}(values[0]);
+{% else %}
+  output[params.offset + i] = values[0];
+{% endif %}
+}

build/webgpu/one-hot-last-axis-vec4.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,117 @@

+{% 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 }}
+// Vec4 OneHot for a contiguous depth-last output. Wide rows are assigned one
+// cooperative workgroup so they share the index load and avoid flat division;
+// short rows use the dense flat launch to avoid underfilled workgroups.
+const DEPTH: u32 = {{ source.depth }}u;
+const VEC_DEPTH: u32 = DEPTH / 4u;
+struct NormalizedIndex {
+  valid: bool,
+  depth: u32,
+};
+fn normalize_index(raw: {{ indexScalar }}) -> NormalizedIndex {
+{% if floatIndices %}
+  let depth_i = i32(DEPTH);
+  let truncated = i32(trunc(raw));
+  let adjusted = select(truncated, truncated + depth_i, truncated < 0);
+  return NormalizedIndex(adjusted >= 0 && adjusted < depth_i, u32(adjusted));
+{% elif usesI32 %}
+  let depth_i = i32(DEPTH);
+  let adjusted = select(raw, raw + depth_i, raw < 0);
+  return NormalizedIndex(adjusted >= 0 && adjusted < depth_i, u32(adjusted));
+{% else %}
+  return NormalizedIndex(raw < DEPTH, raw);
+{% endif %}
+}
+{% if source.cooperative %}
+const ROWS: u32 = {{ source.rows }}u;
+const WG: u32 = {{ source.wg }}u;
+var<workgroup> row_index: {{ indexScalar }};
+@compute @workgroup_size({{ source.wg }})
+fn main(
+  @builtin(local_invocation_id) lid3: vec3<u32>,
+  @builtin(workgroup_id) wid: vec3<u32>,
+  @builtin(num_workgroups) nwg: vec3<u32>
+) {
+  let lane = lid3.x;
+  let row = wid.x + wid.y * nwg.x;
+  if (row >= ROWS) {
+    return;
+  }
+  if (lane == 0u) {
+    row_index = indices[row];
+  }
+  workgroupBarrier();
+  let index = normalize_index(row_index);
+  let off_v = values[0];
+  let on_v = values[1];
+  let row_base = row * VEC_DEPTH;
+  for (var d4 = lane; d4 < VEC_DEPTH; d4 += WG) {
+    let d0 = d4 * 4u;
+    output[row_base + d4] = {{ vectorScalar }}(
+      select(off_v, on_v, index.valid && index.depth == d0),
+      select(off_v, on_v, index.valid && index.depth == d0 + 1u),
+      select(off_v, on_v, index.valid && index.depth == d0 + 2u),
+      select(off_v, on_v, index.valid && index.depth == d0 + 3u)
+    );
+  }
+}
+{% else %}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+{{ flat_index_2d("i4", note="") }}
+  let d4 = i4 % VEC_DEPTH;
+  let d0 = d4 * 4u;
+  let row = i4 / VEC_DEPTH;
+  let index = normalize_index(indices[row]);
+  let off_v = values[0];
+  let on_v = values[1];
+  output[i4] = {{ vectorScalar }}(
+    select(off_v, on_v, index.valid && index.depth == d0),
+    select(off_v, on_v, index.valid && index.depth == d0 + 1u),
+    select(off_v, on_v, index.valid && index.depth == d0 + 2u),
+    select(off_v, on_v, index.valid && index.depth == d0 + 3u)
+  );
+}
+{% endif %}

build/webgpu/one-hot-scatter.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,69 @@

+{% 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 }}
+// Sparse on-value scatter for an output laid out as [outer, depth, inner].
+// One thread owns one input index, so output locations are unique even when
+// several input elements select the same depth value.
+const DEPTH: u32 = {{ source.depth }}u;
+const INNER: u32 = {{ source.inner }}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(note="") }}
+  let raw = indices[i];
+{% if floatIndices %}
+  let depth_i = i32(DEPTH);
+  let truncated = i32(trunc(raw));
+  let adjusted = select(truncated, truncated + depth_i, truncated < 0);
+  let valid = adjusted >= 0 && adjusted < depth_i;
+  let on_depth = u32(adjusted);
+{% elif usesI32 %}
+  let depth_i = i32(DEPTH);
+  let adjusted = select(raw, raw + depth_i, raw < 0);
+  let valid = adjusted >= 0 && adjusted < depth_i;
+  let on_depth = u32(adjusted);
+{% else %}
+  let valid = raw < DEPTH;
+  let on_depth = raw;
+{% endif %}
+  if (valid) {
+    let inner_coord = i % INNER;
+    let outer = i / INNER;
+    output[(outer * DEPTH + on_depth) * INNER + inner_coord] = values[1];
+  }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,790 @@

+{
+  "op": "ai.onnx.OneHot",
+  "cases": [
+    {
+      "name": "float16_indices_depth_and_values",
+      "provenance": {
+        "notes": "Pins the exact ONNX float16 routes for independent indices, depth, and value type variables."
+      },
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "indices": { "dtype": "float16", "shape": [4], "data": { "kind": "values", "values": [-1.0, 0.0, 2.0, 4.0] } },
+        "depth": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [4.0] } },
+        "values": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [-0.5, 2.0] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "int16_indices_depth_and_values",
+      "provenance": {
+        "notes": "Pins widened-storage int16 support across all three independent OneHot type variables."
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "indices": { "dtype": "int16", "shape": [4], "data": { "kind": "values", "values": [-1, 0, 2, 3] } },
+        "depth": { "dtype": "int16", "shape": [], "data": { "kind": "values", "values": [3] } },
+        "values": { "dtype": "int16", "shape": [2], "data": { "kind": "values", "values": [-1234, 2345] } }
+      },
+      "outputs": { "output": { "dtype": "int16", "shape": [3, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "int8_indices_depth_and_values",
+      "provenance": { "notes": "Pins signed 8-bit logical tensors through widened i32 shader storage." },
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "indices": { "dtype": "int8", "shape": [4], "data": { "kind": "values", "values": [-2, -1, 1, 4] } },
+        "depth": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "values": { "dtype": "int8", "shape": [2], "data": { "kind": "values", "values": [-7, 101] } }
+      },
+      "outputs": { "output": { "dtype": "int8", "shape": [4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "uint8_indices_depth_and_values",
+      "provenance": { "notes": "Pins unsigned 8-bit logical tensors through widened u32 shader storage." },
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "indices": { "dtype": "uint8", "shape": [2, 2], "data": { "kind": "values", "values": [0, 3, 4, 1] } },
+        "depth": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "values": { "dtype": "uint8", "shape": [2], "data": { "kind": "values", "values": [2, 251] } }
+      },
+      "outputs": { "output": { "dtype": "uint8", "shape": [2, 4, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "bool_values_and_output",
+      "provenance": { "notes": "Pins ONNX boolean off/on payloads without exposing bool as an index or depth type." },
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "indices": { "dtype": "int32", "shape": [4], "data": { "kind": "values", "values": [-1, 0, 2, 3] } },
+        "depth": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [3] } },
+        "values": { "dtype": "bool", "shape": [2], "data": { "kind": "values", "values": [0, 1] } }
+      },
+      "outputs": { "output": { "dtype": "bool", "shape": [4, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "float32_indices_int16_depth_float32_values",
+      "provenance": {
+        "notes": "An integral float32 index route with independently typed int16 depth connects both standard type variables while the all-int16 sibling retains widened-storage coverage."
+      },
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "indices": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [-1.0, 0.0, 2.0, 3.0] } },
+        "depth": { "dtype": "int16", "shape": [1], "data": { "kind": "values", "values": [3] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [-0.25, 1.5] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "float32_indices_int8_depth_float32_values",
+      "provenance": {
+        "notes": "An integral float32 index route with independently typed int8 depth connects both standard type variables while the all-int8 sibling retains widened-storage coverage."
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "indices": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [-2.0, -1.0, 1.0, 4.0] } },
+        "depth": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [4] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.125, 2.5] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "float32_indices_uint8_depth_float32_values",
+      "provenance": {
+        "notes": "An integral float32 index route with independently typed uint8 depth connects both standard type variables while the all-uint8 sibling retains widened-storage coverage."
+      },
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "indices": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [0.0, 3.0, 4.0, 1.0] } },
+        "depth": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [-1.0, 3.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 4, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "int16_indices_depth_float32_values",
+      "provenance": { "notes": "Pairs int16 indices/depth with the independent float32 value/output route." },
+      "inputs": {
+        "indices": { "dtype": "int16", "shape": [1], "data": { "kind": "values", "values": [-1] } },
+        "depth": { "dtype": "int16", "shape": [], "data": { "kind": "values", "values": [2] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [-0.25, 1.5] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "int8_indices_depth_float32_values",
+      "provenance": { "notes": "Pairs int8 indices/depth with the independent float32 value/output route." },
+      "inputs": {
+        "indices": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "depth": { "dtype": "int8", "shape": [], "data": { "kind": "values", "values": [2] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.125, 2.5] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "uint8_indices_depth_float32_values",
+      "provenance": { "notes": "Pairs uint8 indices/depth with the independent float32 value/output route." },
+      "inputs": {
+        "indices": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } },
+        "depth": { "dtype": "uint8", "shape": [], "data": { "kind": "values", "values": [2] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [-1.0, 3.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2], "tolerance": 0 } }
+    },
+    {
+      "name": "dispatch_cliff_axis0",
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [4], "data": { "kind": "cycle", "values": [0, 1000000, 4194240, 7] } },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [4194241] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4194241, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "rank1_last_axis",
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [0, 2, 1] } },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 4] } }
+    },
+    {
+      "name": "f32_subnormal_values_last_axis_vec4",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.DefaultAxis_float_float_float",
+        "notes": "OneHot payload values may be valid positive subnormal floats; the last-axis vec4 path must not flush on-values to zero."
+      },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [0, 2, 1] } },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1e-40] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "tolerance": 2e-45,
+          "data": {
+            "kind": "values",
+            "values": [9.99994610111476e-41, 0.0, 0.0, 0.0, 0.0, 0.0, 9.99994610111476e-41, 0.0, 0.0, 9.99994610111476e-41, 0.0, 0.0]
+          }
+        }
+      }
+    },
+    {
+      "name": "rank2_f16",
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [2, 2], "data": { "kind": "values", "values": [0, 1, 2, 1] } },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [3] } },
+        "values": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [-1.0, 2.0] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [2, 2, 3] } },
+      "tolerance": 0.001
+    },
+    {
+      "name": "rank2_axis0",
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [2, 2], "data": { "kind": "values", "values": [0, 1, 2, 1] } },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [3] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [-1.0, 2.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 2, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "f32_subnormal_values_axis0_negative_and_oob",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.Axis_Negative_NegIndex_NonDefault_NonZeroOffValue",
+        "notes": "Generic-axis payload edge with subnormal off/on values, negative index normalization, and an out-of-range index that must remain off."
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "indices": { "dtype": "int32", "shape": [4], "data": { "kind": "values", "values": [0, -1, 5, 1] } },
+        "depth": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [3] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1e-40, 2e-40] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "tolerance": 2e-45,
+          "data": {
+            "kind": "values",
+            "values": [1.999989220222952e-40, 9.99994610111476e-41, 9.99994610111476e-41, 9.99994610111476e-41, 9.99994610111476e-41, 9.99994610111476e-41, 9.99994610111476e-41, 1.999989220222952e-40, 9.99994610111476e-41, 1.999989220222952e-40, 9.99994610111476e-41, 9.99994610111476e-41]
+          }
+        }
+      }
+    },
+    {
+      "name": "rank4_indices_last_axis",
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [1, 2, 1, 2], "data": { "kind": "values", "values": [0, 2, 1, 3] } },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [-1.0, 5.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 1, 2, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank5_indices_last_axis",
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [1, 2, 1, 2, 1], "data": { "kind": "values", "values": [0, 2, 1, 3] } },
+        "depth": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [4] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [-1.0, 5.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [1, 2, 1, 2, 1, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank6_indices_axis0_output_rank7",
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "indices": {
+          "dtype": "uint32",
+          "shape": [1, 2, 1, 2, 1, 2],
+          "data": { "kind": "values", "values": [0, 2, 1, 3, 2, 0, 3, 1] }
+        },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [-2.0, 6.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 1, 2, 1, 2, 1, 2], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "scalar_depth",
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [0, 2, 3] } },
+        "depth": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [4] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [-3.0, 7.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 4], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_default_axis_rank2_depth10",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.DefaultAxis"
+      },
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [2, 3], "data": { "kind": "values", "values": [1, 9, 8, 2, 4, 6] } },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 10], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_float_indices_float_depth_default_axis",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.DefaultAxis_float_float_float",
+        "notes": "Covers integral index and depth values carried by ONNX-permitted float32 tensors."
+      },
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "indices": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [1.0, 9.0, 8.0, 2.0, 4.0, 6.0] }
+        },
+        "depth": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [10.0] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 10], "tolerance": 0 } }
+    },
+    {
+      "name": "float32_indices_float16_depth_nonzero_off",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.DefaultAxis_float_float_float_NonZeroOffValue",
+        "notes": "The ORT-derived integral float-index case uses float16 depth to cover the independent I/D type variables; arbitrary off/on values catch accidental zero-fill assumptions."
+      },
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "indices": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [1.0, 9.0, 8.0, 2.0, 4.0, 6.0] }
+        },
+        "depth": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [10.0] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [2.0, 3.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 10], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int32_indices_depth_float_values_default_axis",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.DefaultAxis_int32_float_int32"
+      },
+      "inputs": {
+        "indices": { "dtype": "int32", "shape": [2, 3], "data": { "kind": "values", "values": [1, 9, 8, 2, 4, 6] } },
+        "depth": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 10], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_axis1_rank2_depth10",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.Axis_1"
+      },
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [2, 3], "data": { "kind": "values", "values": [1, 9, 8, 2, 4, 6] } },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 10, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_negative_axis_nondefault_rank2_depth10",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.Axis_Negative_NonDefault"
+      },
+      "attrs": { "axis": -3 },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [2, 3], "data": { "kind": "values", "values": [1, 9, 8, 2, 4, 6] } },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [10, 2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int32_indices_float32_depth_values_nonzero_off_default_axis",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.DefaultAxis_int32_float_int32_NonZeroOffValue",
+        "notes": "The ORT case's depth is adapted from int32 to float32 to cover the independent I/D type variables while preserving its integral value and exact output."
+      },
+      "inputs": {
+        "indices": { "dtype": "int32", "shape": [2, 3], "data": { "kind": "values", "values": [1, 9, 8, 2, 4, 6] } },
+        "depth": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [10.0] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [2.0, 3.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 10], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_default_axis_nonzero_off_value",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.DefaultAxis_NonZeroOffValue"
+      },
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [2, 3], "data": { "kind": "values", "values": [1, 9, 8, 2, 4, 6] } },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [2.0, 3.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 10], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_dim_with_zero_axis0",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.DimWithZero"
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [2, 0], "data": { "kind": "values", "values": [] } },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [10, 2, 0], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_negative_index_axis_negative_nondefault",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.Axis_Negative_NegIndex_NonDefault"
+      },
+      "attrs": { "axis": -3 },
+      "inputs": {
+        "indices": { "dtype": "int32", "shape": [2, 3], "data": { "kind": "values", "values": [1, -1, 8, 2, 4, 6] } },
+        "depth": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [10, 2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_indices_out_of_range_default_axis_nonzero_off",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.DefaultAxis_IndicesOutOfRange_NonZeroOffValue"
+      },
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "indices": { "dtype": "int32", "shape": [2, 3], "data": { "kind": "values", "values": [1, -1, 8, 13, 4, -12] } },
+        "depth": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [2.0, 3.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 10], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_onehot_negative_indices",
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "indices": { "dtype": "int32", "shape": [3], "data": { "kind": "values", "values": [0, -7, -8] } },
+        "depth": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [10] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.0, 3.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 10] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_onehot_negative_indices",
+        "notes": "ONNX int64 metadata/index tensors use framework int32/uint32 slots where representable. Official backend depth is represented with framework uint32 metadata."
+      }
+    },
+    {
+      "name": "onnx_backend_onehot_with_axis",
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "indices": { "dtype": "int32", "shape": [2, 2], "data": { "kind": "values", "values": [1, 9, 2, 4] } },
+        "depth": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [10] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.0, 3.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 10, 2] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_onehot_with_axis",
+        "notes": "Official backend fixture stores integer-looking indices/depth as floats; represented with framework integer metadata tensors. Official backend depth is represented with framework uint32 metadata."
+      }
+    },
+    {
+      "name": "onnx_backend_onehot_with_negative_axis",
+      "attrs": { "axis": -2 },
+      "inputs": {
+        "indices": { "dtype": "int32", "shape": [2, 2], "data": { "kind": "values", "values": [1, 9, 2, 4] } },
+        "depth": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [10] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [1.0, 3.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 10, 2] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_onehot_with_negative_axis",
+        "notes": "Official backend fixture stores integer-looking indices/depth as floats; represented with framework integer metadata tensors. Official backend depth is represented with framework uint32 metadata."
+      }
+    },
+    {
+      "name": "onnx_backend_onehot_without_axis",
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [0, 7, 8] } },
+        "depth": { "dtype": "uint32", "shape": [], "data": { "kind": "values", "values": [12] } },
+        "values": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [2, 5] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [3, 12] } },
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_onehot_without_axis",
+        "notes": "ONNX int64 metadata/index tensors use framework int32/uint32 slots where representable. Official backend depth is represented with framework uint32 metadata."
+      }
+    },
+    {
+      "name": "ort_int32_values_exact_above_float24_default_axis",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.DefaultAxis_int64_int32_float_NonZeroOffValue",
+        "notes": "Uses framework int32 depth instead of ORT's float depth so WebGPU can treat depth as shape metadata."
+      },
+      "inputs": {
+        "indices": { "dtype": "int32", "shape": [4], "data": { "kind": "values", "values": [0, 2, -1, 5] } },
+        "depth": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "values": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [16777217, -16777219] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_int32_values_oob_indices_exact_above_float24",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.DefaultAxis_IndicesOutOfRange_NonZeroOffValue"
+      },
+      "inputs": {
+        "indices": { "dtype": "int32", "shape": [4], "data": { "kind": "values", "values": [1, -1, 4, -7] } },
+        "depth": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "values": { "dtype": "int32", "shape": [2], "data": { "kind": "values", "values": [20000001, -20000003] } }
+      },
+      "outputs": { "output": { "dtype": "int32", "shape": [4, 4], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_axis0_nonzero_off_value_float",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.Axis_0_NonZeroOffValue",
+        "notes": "Pins axis-0 insertion with distinct nonzero off/on values 2 and 3, producing output shape [10,2,3]."
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [2, 3], "data": { "kind": "values", "values": [1, 9, 8, 2, 4, 6] } },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [2.0, 3.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [10, 2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_axis2_nonzero_off_value_float",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.Axis_2_NonZeroOffValue",
+        "notes": "Pins axis-2 insertion with distinct nonzero off/on values 2 and 3, producing output shape [2,3,10]."
+      },
+      "attrs": { "axis": 2 },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [2, 3], "data": { "kind": "values", "values": [1, 9, 8, 2, 4, 6] } },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [2.0, 3.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 10], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_depth_one_all_indices_on",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.DepthOne"
+      },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [0, 0, 0] } },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [1] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [-4.0, 9.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [3, 1], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_axis0_rank2_depth10",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.Axis_0",
+        "notes": "Uses uint32 indices/depth and float32 values in place of ORT's int64 tensors."
+      },
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [2, 3], "data": { "kind": "values", "values": [1, 9, 8, 2, 4, 6] } },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [10, 2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_axis2_rank2_depth10",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.Axis_2",
+        "notes": "Uses uint32 indices/depth and float32 values in place of ORT's int64 tensors."
+      },
+      "attrs": { "axis": 2 },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [2, 3], "data": { "kind": "values", "values": [1, 9, 8, 2, 4, 6] } },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 10], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_axis1_nonzero_off_value_float",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.Axis_1_NonZeroOffValue",
+        "notes": "Uses uint32 indices/depth and float32 values in place of ORT's int64 tensors."
+      },
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [2, 3], "data": { "kind": "values", "values": [1, 9, 8, 2, 4, 6] } },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [2.0, 3.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 10, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_negative_axis_nondefault_nonzero_off_value_float",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.Axis_Negative_NonDefault_NonZeroOffValue",
+        "notes": "Uses uint32 indices/depth and float32 values in place of ORT's int64 tensors."
+      },
+      "attrs": { "axis": -3 },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [2, 3], "data": { "kind": "values", "values": [1, 9, 8, 2, 4, 6] } },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [2.0, 3.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [10, 2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_negative_axis_negative_index_nonzero_off_value",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.Axis_Negative_NegIndex_NonDefault_NonZeroOffValue",
+        "notes": "Uses int32 indices/depth and float32 values in place of ORT's int64 tensors."
+      },
+      "attrs": { "axis": -3 },
+      "inputs": {
+        "indices": { "dtype": "int32", "shape": [2, 3], "data": { "kind": "values", "values": [1, -1, 8, 2, 4, 6] } },
+        "depth": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [2.0, 3.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [10, 2, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_default_axis_indices_out_of_range_zero_off",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.DefaultAxis_IndicesOutOfRange",
+        "notes": "Uses int32 indices/depth and float32 values in place of ORT's int64 tensors."
+      },
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "indices": { "dtype": "int32", "shape": [2, 3], "data": { "kind": "values", "values": [1, -1, 8, 13, 4, -12] } },
+        "depth": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 3, 10], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_default_axis_int32_f16_nonzero_off_value",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/tensor/onehot_op_test.cc",
+        "test": "OneHotOpTest.DefaultAxis_int32_MLFloat16_int32_NonZeroOffValue"
+      },
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "indices": { "dtype": "int32", "shape": [2, 3], "data": { "kind": "values", "values": [1, 9, 8, 2, 4, 6] } },
+        "depth": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [10] } },
+        "values": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [2.0, 3.0] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "float16",
+          "shape": [2, 3, 10],
+          "tolerance": 0,
+          "data": {
+            "kind": "values",
+            "values": [2.0, 3.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 3.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 3.0, 2.0, 2.0, 2.0, 3.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 3.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 2.0, 3.0, 2.0, 2.0, 2.0]
+          }
+        }
+      }
+    },
+    {
+      "name": "f16_last_axis_vec4_depth8",
+      "provenance": {
+        "notes": "f16 OneHot with a depth divisible by 4 selects last_axis_vec4 with usesF16=true (a vec4<f16> 64-bit store). Every existing f16 case has depth 3 or 10 (not %4) and routes to the scalar generic_axis kernel, so the vec4<f16> store branch is entirely untested. A real f16 embedding/classification one-hot over a %4 vocab hits this path."
+      },
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [0, 5, 7] } },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [8] } },
+        "values": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [-1.0, 2.0] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [3, 8], "tolerance": 0.001 } }
+    },
+    {
+      "name": "uint32_values_last_axis_vec4_extremes",
+      "provenance": {
+        "notes": "Standard ONNX uint32 payload coverage. Values above float32's exact-integer range ensure the values/output path remains integer-native while the depth-last vec4 route fills and scatters."
+      },
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [3], "data": { "kind": "values", "values": [0, 3, 7] } },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "values": { "dtype": "uint32", "shape": [2], "data": { "kind": "values", "values": [4294967295, 2147483649] } }
+      },
+      "outputs": {
+        "output": {
+          "dtype": "uint32",
+          "shape": [3, 4],
+          "tolerance": 0,
+          "data": {
+            "kind": "values",
+            "values": [2147483649, 4294967295, 4294967295, 4294967295, 4294967295, 4294967295, 4294967295, 2147483649, 4294967295, 4294967295, 4294967295, 4294967295]
+          }
+        }
+      }
+    },
+    {
+      "name": "float_indices_last_axis_vec4_depth8",
+      "provenance": {
+        "notes": "Float32 indices with a depth divisible by 4 select last_axis_vec4 with floatIndices=true (i32(trunc(raw)) + negative-wrap, then a vec4 store). Every existing float-index case has depth 10 (not %4) and routes to generic_axis, so the floatIndices branch of the vec4 kernel is untested. ONNX permits numeric indices (cast to int); a negative float index exercises the truncation+wrap."
+      },
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "indices": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.0, 6.9, -1.0, 11.0] } },
+        "depth": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [8.0] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "f16_vocab256_default_axis",
+      "provenance": {
+        "notes": "Realistic f16 one-hot over a 256-wide vocab (depth %4) routed through last_axis_vec4 usesF16. Larger %4 depth than the depth8 case so several vec4 lanes are all-off before the on-lane, exercising the vec4 fill + single-lane scatter at a model-realistic width."
+      },
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [5], "data": { "kind": "values", "values": [0, 1, 200, 255, 137] } },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [256] } },
+        "values": { "dtype": "float16", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float16", "shape": [5, 256], "tolerance": 0.001 } }
+    },
+    {
+      "name": "int32_indices_last_axis_vec4_depth8_negwrap",
+      "provenance": {
+        "notes": "int32 indices with depth %4 select last_axis_vec4 usesI32 (negative-wrap + bounds). Covers the vec4 i32 branch with a model-realistic %4 depth and a negative index that must wrap to a valid slot plus an out-of-range index that must stay all-off."
+      },
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "indices": { "dtype": "int32", "shape": [4], "data": { "kind": "values", "values": [0, -1, 3, 99] } },
+        "depth": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [8] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "empty_output_last_axis_vec4_depth8",
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [0], "data": { "kind": "values", "values": [] } },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [8] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [0, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "vec4_uint32_on_slot_last_lane_and_oob_depth8",
+      "attrs": { "axis": -1 },
+      "inputs": {
+        "indices": { "dtype": "uint32", "shape": [4], "data": { "kind": "values", "values": [7, 8, 100, 4] } },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [8] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [-5.0, 9.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 8], "tolerance": 0 } }
+    },
+    {
+      "name": "generic_float_indices_middle_axis_negwrap",
+      "attrs": { "axis": 1 },
+      "inputs": {
+        "indices": {
+          "dtype": "float32",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [0.0, 6.9, -1.0, 3.0, 10.0, 2.5] }
+        },
+        "depth": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [7.0] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [2.0, 3.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [2, 7, 3], "tolerance": 0 } }
+    },
+    {
+      "name": "rank7_indices_axis0_output_rank8",
+      "attrs": { "axis": 0 },
+      "inputs": {
+        "indices": {
+          "dtype": "uint32",
+          "shape": [1, 2, 1, 2, 1, 2, 2],
+          "data": { "kind": "cycle", "values": [0, 3, 1, 2] }
+        },
+        "depth": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [4] } },
+        "values": { "dtype": "float32", "shape": [2], "data": { "kind": "values", "values": [0.0, 1.0] } }
+      },
+      "outputs": { "output": { "dtype": "float32", "shape": [4, 1, 2, 1, 2, 1, 2, 2], "tolerance": 0.000001 } }
+    }
+  ]
+}