sync 91d990483a17

Browse files

Files changed (8) hide show

README.md +12 -8
build/webgpu/bench.json +1 -2
build/webgpu/cast-scalar-x4.wgsl.jinja +2 -5
build/webgpu/manifest.json +66 -87
build/webgpu/metadata.json +16 -9
build/webgpu/test.json +4 -5
build/webgpu/unary-scalar.wgsl.jinja +27 -8
build/webgpu/unary-vec4.wgsl.jinja +30 -8

README.md CHANGED Viewed

@@ -18,16 +18,16 @@ See the [ONNX `CastLike` spec](https://onnx.ai/onnx/operators/onnx__CastLike.htm
 ## 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
@@ -47,7 +47,7 @@ Default values (overridable per request):
 ## 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
@@ -57,10 +57,14 @@ Default values (overridable per request):
 ## 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.

 ## Inputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `x` | `input` | `T1` | — | — | Input tensor whose elements are to be cast. | required |
+| `target` | `target_type` | `T2` | — | — | Tensor whose element type defines the destination dtype; its values are not used. | required |
 ## Outputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `y` | `output` | `T2` | same as `x` | same as `x` | Output tensor with the same shape as `input` and the element type of `target_type`. | required |
 ## Attributes
 ## Files
+- [`metadata.json`](build/webgpu/metadata.json) — kernel metadata (id, digests, per-variant templates, 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
 ## Use with `@huggingface/kernels`
+```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.2
+```
+Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
 The `version: 1` option selects the published kernel contract; it is independent of any operator opset, contrib `since_version`, or model version.
+It follows the `v1` branch as fixes land. To pin exact artifact bytes, pass a 40-character commit `revision` instead of `version`.
 Replace each `*Data` placeholder with a typed array containing the corresponding input data.

build/webgpu/bench.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.CastLike",
   "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
   "cases": [
     {
@@ -53,7 +52,7 @@
       "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] } },

 {
   "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
   "cases": [
     {
       "preset": "smoke",
       "vars": { "inputDtype": "int8", "outputDtype": "float32", "count": 1048579 },
       "inputs": {
+        "x": { "dtype": "int8", "shape": [1048579], "dist": "uniform", "seed": 915, "min": -128, "max": 128 },
         "target": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } }
       },
       "outputs": { "y": { "dtype": "float32", "shape": [1048579] } },

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

@@ -1,6 +1,3 @@
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 {% if wrapNarrowInt %}
@@ -13,11 +10,11 @@ fn cast_wrap(v: f32) -> {{ outScalar }} {
 {% 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) %}

 {{ env.wgsl.resourceDeclarations }}
 {% if wrapNarrowInt %}
 {% endif %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: 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 * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if (invocation != 0u) { return; }
   let base = params.count - params.count % 4u;
 {% for lane in range(4) %}

build/webgpu/manifest.json CHANGED Viewed

@@ -2,87 +2,42 @@
   "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\"))) ~ \">\""
@@ -91,18 +46,25 @@
         {
           "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\"))",
@@ -113,23 +75,35 @@
         {
           "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\""
@@ -138,9 +112,14 @@
         {
           "id": "main",
           "name": "CastLike",
-          "source": { "shader": "unary-scalar.wgsl.jinja", "inputs": { "op": "\"castlike\"" } },
-          "bindings": "scalar",
-          "dispatch": { "threads": "numel(shapes.y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     }

   "domain": "ai.onnx",
   "name": "CastLike",
   "sinceVersion": 25,
+  "inputs": { "x": { "onnx": "input", "dtype": "T1" }, "target": { "onnx": "target_type", "dtype": "T2" } },
+  "outputs": { "y": { "onnx": "output", "dtype": "T2", "rank": "ranks.x", "shape": "shapes.x" } },
+  "attributes": { "round_mode": { "default": "up" }, "saturate": { "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"]
   },
+  "tunables": { "WORKGROUP_SIZE": { "default": 256 }, "VEC4_TAIL_MIN_ELEMENTS": { "default": 4096 } },
+  "derive": {
     "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\""
   },
+  "when": ["numel(shapes.x) == numel(shapes.y)", "ranks.y == ranks.x", "f16Ok(dtypes.T1) and f16Ok(dtypes.T2)"],
+  "bindings": {
+    "x": { "buffer": "read-only-storage", "elementType": "$vectorScalar" },
+    "y": { "buffer": "storage", "elementType": "$outVec" },
+    "params": {
+      "buffer": "uniform",
+      "struct": [{ "name": "count", "type": "u32", "value": "floor(numel(shapes.y) / 4)" }]
+    },
+    "x_2": { "name": "x", "buffer": "read-only-storage", "elementType": "$inScalar" },
+    "y_2": { "name": "y", "buffer": "storage", "elementType": "$outScalar" },
+    "params_2": {
+      "name": "params",
+      "buffer": "uniform",
+      "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }]
+    }
+  },
   "variants": [
     {
       "id": "same_layout_vec4",
+      "priority": 20,
+      "when": ["numel(shapes.x) > 0", "numel(shapes.x) % 4 == 0", "tensorDtypes.y != \"bool\""],
+      "derive": {
         "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\"))) ~ \">\""
         {
           "id": "main",
           "name": "CastLike.vec4",
+          "shader": "unary-vec4.wgsl.jinja",
+          "derive": {
+            "op": "\"castlike\"",
+            "vec4PerThread": "4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1"
+          },
+          "bindings": ["x", "y", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(floor(numel(shapes.y) / 4), 4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(floor(numel(shapes.y) / 4), 4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
+      ]
     },
     {
       "id": "vec4_tail",
       "priority": 10,
+      "when": ["numel(shapes.x) >= tunables.VEC4_TAIL_MIN_ELEMENTS", "numel(shapes.x) % 4 != 0", "tensorDtypes.y != \"bool\""],
+      "derive": {
         "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\"))",
         {
           "id": "bulk",
           "name": "CastLike.vec4Bulk",
+          "shader": "unary-vec4.wgsl.jinja",
+          "derive": {
+            "op": "\"castlike\"",
+            "vec4PerThread": "4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1"
+          },
+          "bindings": ["x", "y", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(floor(numel(shapes.y) / 4), 4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(floor(numel(shapes.y) / 4), 4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         },
         {
           "id": "tail",
           "name": "CastLike.scalarTail",
+          "shader": "cast-scalar-x4.wgsl.jinja",
+          "derive": { "tailOnly": true },
+          "bindings": ["x_2", "y_2", "params_2"],
+          "dispatch": {
+            "x": "min(ceilDiv((1), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((1), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "elementwise",
+      "derive": {
         "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\""
         {
           "id": "main",
           "name": "CastLike",
+          "shader": "unary-scalar.wgsl.jinja",
+          "derive": { "op": "\"castlike\"", "itemsPerInvocation": 4 },
+          "bindings": ["x_2", "y_2", "params_2"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.y), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.y), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     }

build/webgpu/metadata.json CHANGED Viewed

@@ -1,20 +1,27 @@
 {
   "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" }
 }

 {
   "name": "ai.onnx.CastLike",
+  "id": "_ai_onnx_castlike_webgpu_b3b5a5c",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "HvPktrO4Hl0z7C35f4x48CIdfhgr5BXEPqDjvZWDBbw=",
+      "cast-scalar-x4.wgsl.jinja": "ZEmWwTu9NiMzEKa+PuUlxatkSfJ8WLLYD78x0ay8l/U=",
+      "manifest.json": "omTaO2Iw4AaBsEmMGrnaTWCUpxDFQ1mQQtk2ivwV7Yo=",
+      "test.json": "uXxmLeBAgMyAvKw77XQjlYD+MzcMu9YPDObTu9xM9CI=",
+      "unary-scalar.wgsl.jinja": "zBUAwYvX5TCH+YwAb8iuo9hhR7g51kBVeFGOMjxn8oo=",
+      "unary-vec4.wgsl.jinja": "vndf17SxewEvqPOmYVKn4LQTBbVuBNjQ3cDs+olqyTo="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "same_layout_vec4": ["unary-vec4.wgsl.jinja"],
+      "vec4_tail": ["cast-scalar-x4.wgsl.jinja", "unary-vec4.wgsl.jinja"],
+      "elementwise": ["unary-scalar.wgsl.jinja"]
+    }
+  }
 }

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.CastLike",
   "cases": [
     {
       "name": "float32_to_uint32_like",
@@ -17,7 +16,7 @@
       "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": {
@@ -33,7 +32,7 @@
       "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] } },
@@ -281,7 +280,7 @@
       },
       "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."
       }
     },
     {
@@ -342,7 +341,7 @@
       },
       "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."
       }
     }
   ]

 {
   "cases": [
     {
       "name": "float32_to_uint32_like",
       "name": "uint32_to_float32_like_extremes",
       "provenance": {
         "source": "onnx/onnx/docs/Operators.md#CastLike-25",
+        "notes": "Uint32 inputs include values beyond signed-int32 and exact-float32 ranges."
       },
       "inputs": {
         "x": {
       "name": "bool_to_int32_like",
       "provenance": {
         "source": "onnx/onnx/docs/Operators.md#CastLike-25",
+        "notes": "Logical bool values remain distinct from their widened uint32 storage representation during conversion to int32."
       },
       "inputs": {
         "x": { "dtype": "bool", "shape": [4], "data": { "kind": "values", "values": [0, 1, 1, 0] } },
       },
       "outputs": { "y": { "dtype": "uint8", "shape": [4099], "tolerance": 0 } },
       "provenance": {
+        "notes": "Casting 4,099 float16 values to uint8 combines vectorized conversion with a three-element scalar tail. Every source value is exactly representable in float16 and lies within the uint8 range."
       }
     },
     {
       },
       "outputs": { "y": { "dtype": "int32", "shape": [1, 1, 1, 1, 1, 1, 8], "tolerance": 0 } },
       "provenance": {
+        "notes": "A rank-7 flat input exercises float32-to-int32 conversion by flat element count on the shared vec4 path."
       }
     }
   ]

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

@@ -1,18 +1,37 @@
 {% 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 %}

 {% macro flat_tail_open() %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
+  // dispatch's per-axis workgroup fold width (the dispatch caps x and spills the rest into y).
+  let invocation = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
+  // Tail-safe scalar x4 keeps vector-like dispatch density without requiring
+  // the logical tensor length (or its storage binding) to be vec4 aligned.
+{% set itemsPerInvocation = itemsPerInvocation if itemsPerInvocation is defined else 4 %}
+{% if tailOnly is defined and tailOnly %}
+  // Tail of a vec4 bulk pass: lane zero alone covers the elements the packed
+  // pass left behind, from the last multiple of the item count to the end.
+  if (invocation != 0u) {
     return;
+  }
+  let begin = params.count - params.count % {{ itemsPerInvocation }}u;
+  let end = params.count;
+{% else %}
+  let begin = invocation * {{ itemsPerInvocation }}u;
+  let end = min(begin + {{ itemsPerInvocation }}u, params.count);
+{% endif %}
+  for (var i = begin; i < end; i = i + 1u) {
+{%- endmacro %}
 {% macro flat_tail_close() %}
+  }
 {% endmacro %}
+{% set toBool = toBool is defined and toBool %}
+{% set wrapNarrowInt = wrapNarrowInt is defined and wrapNarrowInt %}
+{% set wrapSigned = wrapSigned is defined and wrapSigned %}
+{% set inScalar = inScalar | default(scalar) %}
+{% set outScalar = outScalar | default(scalar) %}
+// Scalar unary elementwise implementation. Specialization emits only the
+// selected operation and any numerical helper it requires.
 {% if usesF16 %}
 enable f16;
 {% endif %}

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

@@ -1,19 +1,38 @@
-// 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:
@@ -23,4 +42,7 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
 {% else %}
   y[i] = {{ outVec }}(xv);
 {% endif %}
 }

+{% set wrapNarrowInt = wrapNarrowInt is defined and wrapNarrowInt %}
+{% set wrapSigned = wrapSigned is defined and wrapSigned %}
+{% set outVec = outVec | default(vectorScalar) %}
+// Loads and stores vec4<T> while evaluating the selected unary operation per
+// component.
 {{ env.wgsl.resourceDeclarations }}
+{% set vec4PerThread = vec4PerThread %}
+{% if vec4PerThread > 1 %}
+const ITEMS: u32 = {{ vec4PerThread }}u;
+{% endif %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits when the element count exceeds the
+  // per-axis dispatch fold width (the dispatch caps x and spills the rest into y).
+{% if vec4PerThread > 1 %}
+  // Each invocation walks ITEMS vec4 groups a span apart. Consecutive lanes
+  // access consecutive words on every step, while each lane can keep several
+  // independent loads in flight.
+  let tid = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
+  let span = (params.count + ITEMS - 1u) / ITEMS;
+  for (var j = 0u; j < ITEMS; j = j + 1u) {
+    let i = tid + j * span;
+    if (i >= params.count) {
+      break;
+    }
+{% else %}
+  let i = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if (i >= params.count) {
     return;
   }
+{% endif %}
   let xv = x[i];
 {% if wrapNarrowInt %}
   // Float-to-int8/uint8 Cast and CastLike wrap modulo 256 rather than saturate:
 {% else %}
   y[i] = {{ outVec }}(xv);
 {% endif %}
+{% if vec4PerThread > 1 %}
+  }
+{% endif %}
 }