sync 2e7068faf55e

Browse files

Files changed (8) hide show

README.md +69 -0
build/webgpu/bench.json +82 -0
build/webgpu/cast-scalar-x4.wgsl.jinja +28 -0
build/webgpu/manifest.json +148 -0
build/webgpu/metadata.json +20 -0
build/webgpu/test.json +349 -0
build/webgpu/unary-scalar.wgsl.jinja +39 -0
build/webgpu/unary-vec4.wgsl.jinja +26 -0

README.md CHANGED Viewed

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

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.CastLike
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 25
+## Description
+Casts every element of `input` to the same dtype as `target_type`, producing an output with the same shape as `input`. The `target_type` tensor itself is used only for its dtype and is not read elementwise.
+See the [ONNX `CastLike` spec](https://onnx.ai/onnx/operators/onnx__CastLike.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `input` | `x` | `T1` | — | — | Input tensor whose elements are to be cast. | required |
+| `target_type` | `target` | `T2` | — | — | Tensor whose element type defines the destination dtype; its values are not used. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `output` | `y` | `T2` | same as `input` | same as `input` | Output tensor with the same shape as `input` and the element type of `target_type`. | required |
+## Attributes
+Default values (overridable per request):
+| Attribute | Default | Description |
+| --- | --- | --- |
+| `round_mode` | `"up"` | Rounding direction used only when casting to float8e8m0. The implemented non-float8 subset accepts the ONNX default `"up"`. |
+| `saturate` | `1` | Whether casts to float8 saturate at the finite range. The implemented non-float8 subset accepts the ONNX default `1`. |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T1` | `float32`, `float16`, `uint32`, `int32`, `uint8`, `int8`, `bool` |
+| `T2` | `float32`, `float16`, `uint32`, `int32`, `uint8`, `int8`, `bool` |
+## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, provenance)
+- [`manifest.json`](build/webgpu/manifest.json) — the op contract (source of truth)
+- [`test.json`](build/webgpu/test.json) — correctness cases
+- [`bench.json`](build/webgpu/bench.json) — benchmark + tuning cases
+- [`cast-scalar-x4.wgsl.jinja`](build/webgpu/cast-scalar-x4.wgsl.jinja)
+- [`unary-scalar.wgsl.jinja`](build/webgpu/unary-scalar.wgsl.jinja)
+- [`unary-vec4.wgsl.jinja`](build/webgpu/unary-vec4.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.CastLike", { version: 1 });
+const { y } = await kernel({ x: { data: xData, shape: [] }, target: { data: targetData, shape: [3] } });
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,82 @@

+{
+  "op": "ai.onnx.CastLike",
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "castlike-f32-to-u32-1m",
+      "preset": "smoke",
+      "vars": { "inputDtype": "float32", "outputDtype": "uint32", "count": 1048576 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1048576], "dist": "uniform", "seed": 912, "scale": 255 },
+        "target": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "uint32", "shape": [1048576] } },
+      "bench": {
+        "primary": true,
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "castlike-f32-to-u32-1m-plus3",
+      "preset": "smoke",
+      "vars": { "inputDtype": "float32", "outputDtype": "uint32", "count": 1048579 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1048579], "dist": "uniform", "seed": 913, "scale": 255 },
+        "target": { "dtype": "uint32", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "y": { "dtype": "uint32", "shape": [1048579] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "castlike-f32-to-f16-1m",
+      "preset": "smoke",
+      "vars": { "inputDtype": "float32", "outputDtype": "float16", "count": 1048576 },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1048576], "dist": "normal", "seed": 914 },
+        "target": { "dtype": "float16", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [1048576] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "castlike-i8-to-f32-1m-plus3",
+      "preset": "smoke",
+      "vars": { "inputDtype": "int8", "outputDtype": "float32", "count": 1048579 },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [1048579], "dist": "uniform", "seed": 915, "scale": 127, "signed": true },
+        "target": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1048579] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    },
+    {
+      "name": "castlike-f16-to-f32-unaligned-1m-plus1",
+      "preset": "edge",
+      "vars": { "inputDtype": "float16", "outputDtype": "float32", "count": 1048577 },
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [1048577], "dist": "normal", "seed": 9231 },
+        "target": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1048577] } },
+      "bench": {
+        "metrics": [
+          { "type": "bandwidth", "value": "args.count * (dtypeBytes(args.inputDtype) + dtypeBytes(args.outputDtype))" }
+        ]
+      }
+    }
+  ]
+}

build/webgpu/cast-scalar-x4.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,28 @@

+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+{% if wrapNarrowInt %}
+// ONNX float-to-int8/uint8 narrowing wraps modulo 256 rather than saturating.
+// Preserve the low byte after truncation; the signed path then sign-extends it.
+fn cast_wrap(v: f32) -> {{ outScalar }} {
+  let low = i32(v) & 0xFF;
+  return {{ outScalar }}({% if wrapSigned %}select(low, low - 256, low > 127){% else %}low{% endif %});
+}
+{% endif %}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+  // Four scalar conversions per invocation retain vector-like dispatch density even when the
+  // logical element count is not vec4 aligned. For a vec4 bulk + scalar tail plan, only lane zero
+  // of the tail dispatch runs and starts at the first element not covered by the packed pass.
+  let invocation = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if (invocation != 0u) { return; }
+  let base = params.count - params.count % 4u;
+{% for lane in range(4) %}
+  if (base + {{ lane }}u < params.count) {
+    y[base + {{ lane }}u] = {% if wrapNarrowInt %}cast_wrap(f32(x[base + {{ lane }}u])){% else %}{{ outScalar }}(x[base + {{ lane }}u]){% endif %};
+  }
+{% endfor %}
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,148 @@

+{
+  "domain": "ai.onnx",
+  "name": "CastLike",
+  "sinceVersion": 25,
+  "description": "Casts every element of `input` to the same dtype as `target_type`, producing an output with the same shape as `input`. The `target_type` tensor itself is used only for its dtype and is not read elementwise.",
+  "inputs": [
+    { "role": "input", "dtype": "T1", "description": "Input tensor whose elements are to be cast." },
+    {
+      "role": "target_type",
+      "dtype": "T2",
+      "description": "Tensor whose element type defines the destination dtype; its values are not used."
+    }
+  ],
+  "outputs": [
+    {
+      "role": "output",
+      "dtype": "T2",
+      "rank": "ranks.input",
+      "description": "Output tensor with the same shape as `input` and the element type of `target_type`.",
+      "shape": "shapes.input"
+    }
+  ],
+  "attributes": { "round_mode": "up", "saturate": 1 },
+  "attributeDescriptions": {
+    "round_mode": "Rounding direction used only when casting to float8e8m0. The implemented non-float8 subset accepts the ONNX default `\"up\"`.",
+    "saturate": "Whether casts to float8 saturate at the finite range. The implemented non-float8 subset accepts the ONNX default `1`."
+  },
+  "attributeConstraints": { "round_mode": { "values": ["up"] }, "saturate": { "values": [1] } },
+  "typeConstraints": {
+    "T1": ["float32", "float16", "uint32", "int32", "uint8", "int8", "bool"],
+    "T2": ["float32", "float16", "uint32", "int32", "uint8", "int8", "bool"]
+  },
+  "args": {
+    "x": { "kind": "tensor", "semantic": "input", "role": "input" },
+    "target": { "kind": "tensor", "semantic": "target_type", "role": "input" },
+    "y": { "kind": "tensor", "semantic": "output", "role": "output" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256, "VEC4_TAIL_MIN_ELEMENTS": 4096 },
+  "bindingSets": {
+    "vec4": [
+      {
+        "name": "x",
+        "arg": "x",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$vectorScalar"
+      },
+      { "name": "y", "arg": "y", "semantic": "output", "buffer": { "type": "storage" }, "elementType": "$outVec" },
+      {
+        "name": "params",
+        "semantic": "kernel.params",
+        "buffer": { "type": "uniform" },
+        "struct": {
+          "name": "Params",
+          "fields": [{ "name": "count", "type": "u32", "value": "floor(numel(shapes.y) / 4)" }]
+        }
+      }
+    ],
+    "scalar": [
+      {
+        "name": "x",
+        "arg": "x",
+        "semantic": "input",
+        "buffer": { "type": "read-only-storage" },
+        "elementType": "$inScalar"
+      },
+      { "name": "y", "arg": "y", "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.y)" }] }
+      }
+    ]
+  },
+  "constants": {
+    "usesF16": "tensorDtypes.x == \"float16\" or tensorDtypes.y == \"float16\"",
+    "wrapNarrowInt": "(tensorDtypes.x == \"float32\" or tensorDtypes.x == \"float16\") and (tensorDtypes.y == \"int8\" or tensorDtypes.y == \"uint8\")",
+    "wrapSigned": "tensorDtypes.y == \"int8\""
+  },
+  "variants": [
+    {
+      "id": "same_layout_vec4",
+      "when": ["numel(shapes.x) > 0", "numel(shapes.x) % 4 == 0", "numel(shapes.x) == numel(shapes.y)", "ranks.y == ranks.x", "tensorDtypes.y != \"bool\"", "f16Ok(dtypes.T1) and f16Ok(dtypes.T2)"],
+      "constants": {
+        "scalar": "\"f16\" if tensorDtypes.x == \"float16\" else (\"f32\" if tensorDtypes.x == \"float32\" else (\"i32\" if tensorDtypes.x == \"int32\" or tensorDtypes.x == \"int8\" else \"u32\"))",
+        "vectorScalar": "\"vec4<\" ~ (\"f16\" if tensorDtypes.x == \"float16\" else (\"f32\" if tensorDtypes.x == \"float32\" else (\"i32\" if tensorDtypes.x == \"int32\" or tensorDtypes.x == \"int8\" else \"u32\"))) ~ \">\"",
+        "outVec": "\"vec4<\" ~ (\"f16\" if tensorDtypes.y == \"float16\" else (\"f32\" if tensorDtypes.y == \"float32\" else (\"i32\" if tensorDtypes.y == \"int32\" or tensorDtypes.y == \"int8\" else \"u32\"))) ~ \">\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "CastLike.vec4",
+          "source": { "shader": "unary-vec4.wgsl.jinja", "inputs": { "op": "\"castlike\"" } },
+          "bindings": "vec4",
+          "dispatch": { "threads": "floor(numel(shapes.y) / 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ],
+      "priority": 20
+    },
+    {
+      "id": "vec4_tail",
+      "priority": 10,
+      "when": ["numel(shapes.x) >= tunables.VEC4_TAIL_MIN_ELEMENTS", "numel(shapes.x) % 4 != 0", "ranks.y == ranks.x", "numel(shapes.x) == numel(shapes.y)", "tensorDtypes.y != \"bool\"", "f16Ok(dtypes.T1) and f16Ok(dtypes.T2)"],
+      "constants": {
+        "inScalar": "\"f16\" if tensorDtypes.x == \"float16\" else (\"f32\" if tensorDtypes.x == \"float32\" else (\"i32\" if tensorDtypes.x == \"int32\" or tensorDtypes.x == \"int8\" else \"u32\"))",
+        "outScalar": "\"f16\" if tensorDtypes.y == \"float16\" else (\"f32\" if tensorDtypes.y == \"float32\" else (\"i32\" if tensorDtypes.y == \"int32\" or tensorDtypes.y == \"int8\" else \"u32\"))",
+        "scalar": "\"f16\" if tensorDtypes.x == \"float16\" else (\"f32\" if tensorDtypes.x == \"float32\" else (\"i32\" if tensorDtypes.x == \"int32\" or tensorDtypes.x == \"int8\" else \"u32\"))",
+        "vectorScalar": "\"vec4<\" ~ (\"f16\" if tensorDtypes.x == \"float16\" else (\"f32\" if tensorDtypes.x == \"float32\" else (\"i32\" if tensorDtypes.x == \"int32\" or tensorDtypes.x == \"int8\" else \"u32\"))) ~ \">\"",
+        "outVec": "\"vec4<\" ~ (\"f16\" if tensorDtypes.y == \"float16\" else (\"f32\" if tensorDtypes.y == \"float32\" else (\"i32\" if tensorDtypes.y == \"int32\" or tensorDtypes.y == \"int8\" else \"u32\"))) ~ \">\""
+      },
+      "passes": [
+        {
+          "id": "bulk",
+          "name": "CastLike.vec4Bulk",
+          "source": { "shader": "unary-vec4.wgsl.jinja", "inputs": { "op": "\"castlike\"" } },
+          "bindings": "vec4",
+          "dispatch": { "threads": "floor(numel(shapes.y) / 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        },
+        {
+          "id": "tail",
+          "name": "CastLike.scalarTail",
+          "source": { "shader": "cast-scalar-x4.wgsl.jinja", "inputs": { "tailOnly": true } },
+          "bindings": "scalar",
+          "dispatch": { "threads": 1, "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    },
+    {
+      "id": "elementwise",
+      "when": ["ranks.y == ranks.x", "numel(shapes.x) == numel(shapes.y)", "f16Ok(dtypes.T1) and f16Ok(dtypes.T2)"],
+      "constants": {
+        "inScalar": "\"f16\" if tensorDtypes.x == \"float16\" else (\"f32\" if tensorDtypes.x == \"float32\" else (\"i32\" if tensorDtypes.x == \"int32\" or tensorDtypes.x == \"int8\" else \"u32\"))",
+        "outScalar": "\"f16\" if tensorDtypes.y == \"float16\" else (\"f32\" if tensorDtypes.y == \"float32\" else (\"i32\" if tensorDtypes.y == \"int32\" or tensorDtypes.y == \"int8\" else \"u32\"))",
+        "toBool": "tensorDtypes.y == \"bool\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "CastLike",
+          "source": { "shader": "unary-scalar.wgsl.jinja", "inputs": { "op": "\"castlike\"" } },
+          "bindings": "scalar",
+          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

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

+{
+  "name": "ai.onnx.CastLike",
+  "id": "_ai_onnx_castlike_webgpu_ff0741d",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "LJNtd0rYJjdKViAHum941BgDF6evzLmd2ZI+y2rZKLo=",
+      "cast-scalar-x4.wgsl.jinja": "PrFeLTvZUbu1qrf0L9BWXSct/kRBmcocjA92lM+hgZw=",
+      "manifest.json": "wafZ0zJmjB3QtZJo7odf1U/4kyIu84s59gKS/h1ItRo=",
+      "test.json": "t27WIYD7Q5vGZKbB3Z9vswIfpPiC3XMvUQyzhSfdAIE=",
+      "unary-scalar.wgsl.jinja": "tIRx4X0JJnKcI45N/xWT7E1rwi2vm7nuKvqNE4eFwaE=",
+      "unary-vec4.wgsl.jinja": "Q77arupT4nkbLU7bJ1Gm43gAtXqjH3roPVo7q7g2txA="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.CastLike" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,349 @@

+{
+  "op": "ai.onnx.CastLike",
+  "cases": [
+    {
+      "name": "float32_to_uint32_like",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [0.1, 1.9, 2.2, 5.8, 7.0, 9.6] }
+        },
+        "target": { "dtype": "uint32", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "uint32", "shape": [6] } }
+    },
+    {
+      "name": "uint32_to_float32_like_extremes",
+      "provenance": {
+        "source": "onnx/onnx/docs/Operators.md#CastLike-25",
+        "notes": "Provides executable source-side uint32 coverage, including values beyond signed-i32 and exact-f32 ranges."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "uint32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [0, 16777217, 2147483648, 4294967295] }
+        },
+        "target": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4], "tolerance": 0 } }
+    },
+    {
+      "name": "bool_to_int32_like",
+      "provenance": {
+        "source": "onnx/onnx/docs/Operators.md#CastLike-25",
+        "notes": "Provides executable source-side bool coverage while keeping the logical bool contract distinct from widened u32 storage."
+      },
+      "inputs": {
+        "x": { "dtype": "bool", "shape": [4], "data": { "kind": "values", "values": [0, 1, 1, 0] } },
+        "target": { "dtype": "int32", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [4], "tolerance": 0 } }
+    },
+    {
+      "name": "int32_to_float32_like_scalar_target",
+      "inputs": {
+        "x": { "dtype": "int32", "shape": [2, 3], "data": { "kind": "values", "values": [-3, -1, 0, 1, 7, 12] } },
+        "target": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [2, 3] } }
+    },
+    {
+      "name": "uint8_to_float16_like",
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [0, 1, 127, 255] } },
+        "target": { "dtype": "float16", "shape": [2, 2], "data": { "kind": "values", "values": [0.0, 0.0, 0.0, 0.0] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [4] } },
+      "tolerance": 0.001
+    },
+    {
+      "name": "float16_to_int32_like",
+      "inputs": {
+        "x": { "dtype": "float16", "shape": [5], "data": { "kind": "values", "values": [-2.75, -1.1, 0.0, 1.9, 4.5] } },
+        "target": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [5] } }
+    },
+    {
+      "name": "ort_function_scalar_float_to_ranked_float_like",
+      "provenance": {
+        "source": "onnxruntime/test/framework/function_test.cc",
+        "test": "FunctionTest.AttrWithDefault",
+        "notes": "Covers ORT's use of CastLike inside a function body; the target tensor contributes only the output dtype, not the output shape."
+      },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [2.0] } },
+        "target": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [0.0, 0.0, 0.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_scatternd_float32_to_float16_like_initializer",
+      "provenance": {
+        "source": "onnxruntime/test/python/onnxruntime_test_scatternd.py",
+        "test": "TestScatterND.common_scatter",
+        "notes": "Adapts ORT's ScatterND helper pattern where CastLike converts float input to the dtype of a float16 initializer before later graph ops."
+      },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [2, 2], "data": { "kind": "values", "values": [-3.5, -0.25, 1.5, 8.0] } },
+        "target": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 2], "tolerance": 0.001 } }
+    },
+    {
+      "name": "onnx_backend_empty_like_float32_to_int8",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/case/node/castlike.py",
+        "test": "CastLike.export",
+        "notes": "Uses the ONNX backend generator pattern where the like tensor is empty and only supplies the target dtype."
+      },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [5], "data": { "kind": "values", "values": [-3.9, -1.1, 0.0, 1.9, 127.9] } },
+        "target": { "dtype": "int8", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [5], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_empty_like_float32_to_uint8",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/case/node/castlike.py",
+        "test": "CastLike.export",
+        "notes": "Uses the ONNX backend generator pattern where the like tensor is empty and only supplies the target dtype."
+      },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [0.9, 1.9, 127.9, 255.9] } },
+        "target": { "dtype": "uint8", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [4], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_empty_like_int8_to_float32",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/case/node/castlike.py",
+        "test": "CastLike.export",
+        "notes": "Uses the ONNX backend generator pattern where the like tensor is empty and only supplies the target dtype."
+      },
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [4], "data": { "kind": "values", "values": [-128, -1, 0, 127] } },
+        "target": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_empty_like_uint8_to_int32",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/case/node/castlike.py",
+        "test": "CastLike.export",
+        "notes": "Uses the ONNX backend generator pattern where the like tensor is empty and only supplies the target dtype."
+      },
+      "inputs": {
+        "x": { "dtype": "uint8", "shape": [4], "data": { "kind": "values", "values": [0, 1, 127, 255] } },
+        "target": { "dtype": "int32", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [4], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_castlike_float_to_float16",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_castlike_FLOAT_to_FLOAT16",
+        "test": "test_castlike_FLOAT_to_FLOAT16"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.4789254665374756, 0.48033666610717773, 0.4996848702430725, 0.8191054463386536, 0.4703124761581421, 0.8164680004119873, 0.21087194979190826, 0.7229037880897522, "NaN", "Infinity", "Infinity", "-Infinity"]
+          }
+        },
+        "target": { "dtype": "float16", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [3, 4], "tolerance": 0.001, "allowNaN": true } }
+    },
+    {
+      "name": "onnx_backend_castlike_float16_to_float",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_castlike_FLOAT16_to_FLOAT",
+        "test": "test_castlike_FLOAT16_to_FLOAT"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [3, 4],
+          "data": {
+            "kind": "values",
+            "values": [0.47900390625, 0.480224609375, 0.499755859375, 0.8193359375, 0.47021484375, 0.81640625, 0.2108154296875, 0.72314453125, "NaN", "Infinity", "Infinity", "-Infinity"]
+          }
+        },
+        "target": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4], "tolerance": 0, "allowNaN": true } }
+    },
+    {
+      "name": "vec4_i32_to_f32_lanes",
+      "inputs": {
+        "x": {
+          "dtype": "int32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [-100, -3, -1, 0, 1, 7, 12, 100] }
+        },
+        "target": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [8], "tolerance": 0 } }
+    },
+    {
+      "name": "vec4_f32_to_i32_like_truncates_toward_zero",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [1.0, 2.9, -3.0, 0.0, 100.0, -0.5, -2.9, 127.75] }
+        },
+        "target": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [8], "tolerance": 0 } }
+    },
+    {
+      "name": "vec4_f16_to_i32_like_truncates_toward_zero",
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "values", "values": [1.0, 2.5, -3.0, 0.0, 100.0, -0.5, -2.5, 7.5] }
+        },
+        "target": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [8], "tolerance": 0 } }
+    },
+    {
+      "name": "vec4_f32_to_i8_like_in_range_truncates",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [-128.0, -1.9, -0.5, 0.0, 1.9, 127.0, -127.75, 126.5] }
+        },
+        "target": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [8], "tolerance": 0 } }
+    },
+    {
+      "name": "scalar_x4_f32_to_u32_like_tail",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [17],
+          "data": {
+            "kind": "values",
+            "values": [0.0, 1.9, 2.2, 3.8, 4.0, 5.6, 6.1, 7.9, 8.0, 9.2, 10.7, 11.0, 12.4, 13.8, 14.0, 15.9, 16.2]
+          }
+        },
+        "target": { "dtype": "uint32", "shape": [0], "data": { "kind": "values", "values": [] } }
+      },
+      "outputs": { "y": { "dtype": "uint32", "shape": [17], "tolerance": 0 } }
+    },
+    {
+      "name": "vec4_tail_i8_to_f32_like_4097",
+      "inputs": {
+        "x": { "dtype": "int8", "shape": [4097], "data": { "kind": "constant", "value": -3 } },
+        "target": { "dtype": "float32", "shape": [1], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4097], "tolerance": 0 } },
+      "provenance": { "notes": "Exercises a packed vec4 bulk followed by a scalar tail in the same dispatch." }
+    },
+    {
+      "name": "vec4_tail_f32_to_i8_like_4099",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4099],
+          "data": { "kind": "cycle", "values": [1.5, -2.5, 126.25, -127.75, 63.5] }
+        },
+        "target": { "dtype": "int8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [4099], "tolerance": 0 } },
+      "provenance": {
+        "notes": "CastLike to an int8 target above the vec4-tail crossover (4099 = 4096 packed elements plus a 3-element scalar tail), selecting the vec4 bulk plus scalar-tail path with its signed int8 storage conversion. Every value is within int8's representable range, so the fixture asserts ONNX-defined truncation toward zero without depending on undefined out-of-range float-to-fixed-point behavior."
+      }
+    },
+    {
+      "name": "vec4_tail_f16_to_u8_like_4099",
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [4099],
+          "data": { "kind": "cycle", "values": [0.5, 200.75, 255.0, 254.25, 1.5, 44.5, 127.5] }
+        },
+        "target": { "dtype": "uint8", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "uint8", "shape": [4099], "tolerance": 0 } },
+      "provenance": {
+        "notes": "Float16 source with a uint8 target above the vec4-tail crossover: renders the scalar tail kernel with usesF16 and wrapNarrowInt both set and the UNSIGNED low-byte return. Every cycle value is exactly representable in float16 and lies within uint8's representable range. ONNX defines out-of-range float-to-fixed-point conversion as undefined, so this portable fixture deliberately does not assert wraparound for such values."
+      }
+    },
+    {
+      "name": "int32_to_bool_like",
+      "inputs": {
+        "x": { "dtype": "int32", "shape": [2, 3], "data": { "kind": "values", "values": [-3, 1, 0, 42, -1, 0] } },
+        "target": { "dtype": "bool", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "bool",
+          "shape": [2, 3],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [1, 1, 0, 1, 1, 0] }
+        }
+      }
+    },
+    {
+      "name": "vec4_f32_finite_overflow_to_float16_like_saturate_vs_inf",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": [70000.0, -1e+30, 65504.0, 65600.0, 1.0, -2.5, 0.0, -70000.0] }
+        },
+        "target": { "dtype": "float16", "shape": [1], "data": { "kind": "values", "values": [0.0] } }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [8], "tolerance": 0.001 } }
+    },
+    {
+      "name": "vec4_f32_nonzero_special_to_bool_like",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": ["NaN", "Infinity", "-Infinity", 0.0, 0.0, 1.0, -2.5, 1e-30] }
+        },
+        "target": { "dtype": "bool", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "bool",
+          "shape": [8],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [1, 1, 1, 0, 0, 1, 1, 1] }
+        }
+      }
+    },
+    {
+      "name": "rank7_vec4_f32_to_i32",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1, 1, 1, 1, 1, 1, 8],
+          "data": { "kind": "values", "values": [0.1, 1.9, 2.2, 5.8, 7.0, 9.6, -3.7, -1.2] }
+        },
+        "target": { "dtype": "int32", "shape": [1], "data": { "kind": "values", "values": [0] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1, 1, 1, 1, 8], "tolerance": 0 } },
+      "provenance": {
+        "notes": "Rank-7 flat input on the shared vec4 unary cast kernel (indexes by flat numel; sibling Cast has no rank cap). Covers the regime the removed 'ranks.x <= 6' guard rejected."
+      }
+    }
+  ]
+}

build/webgpu/unary-scalar.wgsl.jinja ADDED Viewed

	@@ -0,0 +1,39 @@

+{% macro flat_tail_open() %}
+@compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
+  // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
+  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills the rest into y).
+  let invocation = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  let i = invocation;
+  if (i >= params.count) {
+    return;
+  }{% endmacro %}
+{% macro flat_tail_close() %}
+{% endmacro %}
+// Scalar unary fallback. Each branch retains the operation's numeric hardening,
+// including Payne-Hanek trigonometric range reduction and NaN/overflow guards.
+{% if usesF16 %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+{{ flat_tail_open() }}
+{% if toBool %}
+{% if inScalar == "f16" or inScalar == "f32" %}
+  // Bool cast on the raw bits with the sign cleared: NaN != 0 must be true,
+  // but no-NaN fast math can fold the ordered float comparison. Both zeros
+  // clear to 0; everything
+  // else — NaN, infinities, subnormals — is nonzero.
+  y[i] = select(0u, 1u, (bitcast<u32>(f32(x[i])) & 0x7fffffffu) != 0u);
+{% else %}
+  y[i] = select(0u, 1u, x[i] != {{ inScalar }}(0));
+{% endif %}
+{% elif wrapNarrowInt %}
+  // Float-to-int8/uint8 CastLike wraps modulo 256 rather than saturating.
+  let low = i32(f32(x[i])) & 0xFF;
+  y[i] = {{ outScalar }}({% if wrapSigned %}select(low, low - 256, low > 127){% else %}low{% endif %});
+{% else %}
+  y[i] = {{ outScalar }}(x[i]);
+{% endif %}
+{{ flat_tail_close() -}}
+}

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

	@@ -0,0 +1,26 @@

+// Loads and stores vec4<T> (128 bits) while retaining scalar per-component
+// arithmetic, including per-component helper calls for guard-heavy operations.
+{% 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>) {
+  // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
+  // maxComputeWorkgroupsPerDimension limit (the dispatch caps x and spills the rest into y).
+  let i = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
+  if (i >= params.count) {
+    return;
+  }
+  let xv = x[i];
+{% if wrapNarrowInt %}
+  // Float-to-int8/uint8 Cast and CastLike wrap modulo 256 rather than saturate:
+  // retain the low byte of the truncated i32, then sign-extend int8.
+  let low = vec4<i32>(vec4<f32>(xv)) & vec4<i32>(0xFF);
+  y[i] = {{ outVec }}({% if wrapSigned %}select(low, low - vec4<i32>(256), low > vec4<i32>(127)){% else %}low{% endif %});
+{% else %}
+  y[i] = {{ outVec }}(xv);
+{% endif %}
+}