sync 91d990483a17

Browse files

Files changed (7) hide show

README.md +15 -12
build/webgpu/bench.json +0 -1
build/webgpu/datamove-elementwise-copy.wgsl.jinja +0 -3
build/webgpu/datamove-flat-copy.wgsl.jinja +1 -4
build/webgpu/manifest.json +27 -80
build/webgpu/metadata.json +11 -8
build/webgpu/test.json +3 -4

README.md CHANGED Viewed

@@ -18,16 +18,16 @@ See the [ONNX `Squeeze` spec](https://onnx.ai/onnx/operators/onnx__Squeeze.html)
 ## Inputs
-| Name | Bind key | Logical dtype | WebGPU storage | Rank | Shape | Description | Presence |
-| --- | --- | --- | --- | --- | --- | --- | --- |
-| `data` | `data` | `T` | runtime-selected; narrow integers and bool use 32-bit slots | — | — | Input tensor with at least as many dimensions as the largest axis index specified. | required |
-| `axes` | `axes` | `S` | `int32` | `1` | — | Logical int64 1-D tensor of axis indices to squeeze; negative values count from the back, signed axes use int32 WebGPU storage, and omission squeezes all size-1 dimensions. | optional |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
-| --- | --- | --- | --- | --- | --- | --- |
-| `squeezed` | `squeezed` | `T` | — | — | Reshaped tensor with the same data as the input but with the specified size-1 dimensions removed. | required |
 ## Type constraints
@@ -38,7 +38,7 @@ See the [ONNX `Squeeze` spec](https://onnx.ai/onnx/operators/onnx__Squeeze.html)
 ## 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
@@ -47,15 +47,18 @@ See the [ONNX `Squeeze` spec](https://onnx.ai/onnx/operators/onnx__Squeeze.html)
 ## 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:
-- `squeezed`
-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.

 ## Inputs
+| Name | Logical dtype | WebGPU storage | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `data` | `T` | runtime-selected; narrow integers and bool use 32-bit slots | — | — | Input tensor with at least as many dimensions as the largest axis index specified. | required |
+| `axes` | `S` | `int32` | `1` | — | Logical int64 1-D tensor of axis indices to squeeze; negative values count from the back, signed axes use int32 WebGPU storage, and omission squeezes all size-1 dimensions. | optional |
 ## Outputs
+| Name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `squeezed` | `T` | — | — | Reshaped tensor with the same data as the input but with the specified size-1 dimensions removed. | required |
 ## Type constraints
 ## 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
+```
+Outputs with inferable metadata are allocated automatically. Explicit `outputs` entries request optional results or provide metadata that cannot be inferred from the supplied inputs and attributes.
+This example supplies explicit metadata for:
+- `squeezed`
 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.Squeeze",
   "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
   "cases": [
     {

 {
   "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
   "cases": [
     {

build/webgpu/datamove-elementwise-copy.wgsl.jinja CHANGED Viewed

@@ -1,6 +1,3 @@
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;





1	{{ env.wgsl.resourceDeclarations }}
2
3	const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;

build/webgpu/datamove-flat-copy.wgsl.jinja CHANGED Viewed

@@ -1,9 +1,6 @@
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
-const COUNT: u32 = {{ source.count }}u;
 const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})

 {{ env.wgsl.resourceDeclarations }}
+const COUNT: u32 = {{ count }}u;
 const WG: u32 = {{ tunables.WORKGROUP_SIZE }}u;
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})

build/webgpu/manifest.json CHANGED Viewed

@@ -2,114 +2,61 @@
   "domain": "ai.onnx",
   "name": "Squeeze",
   "sinceVersion": 13,
-  "description": "Removes single-dimensional entries from the shape of a tensor. If `axes` is provided, only those dimensions are removed; if omitted, all size-1 dimensions are removed. Selecting an axis whose size is not 1 raises an error.",
-  "inputs": [
-    {
-      "role": "data",
-      "dtype": "T",
-      "description": "Input tensor with at least as many dimensions as the largest axis index specified."
-    },
-    {
-      "role": "axes",
-      "dtype": "S",
-      "rank": 1,
-      "optional": true,
-      "description": "Logical int64 1-D tensor of axis indices to squeeze; negative values count from the back, signed axes use int32 WebGPU storage, and omission squeezes all size-1 dimensions."
-    }
-  ],
-  "outputs": [
-    {
-      "role": "squeezed",
-      "dtype": "T",
-      "description": "Reshaped tensor with the same data as the input but with the specified size-1 dimensions removed."
-    }
-  ],
   "typeConstraints": {
     "T": ["float32", "float16", "int32", "uint32", "int16", "uint8", "int8", "bool"],
     "S": ["int64"]
   },
-  "args": {
-    "data": { "kind": "tensor", "semantic": "data", "role": "input" },
-    "axes": {
-      "kind": "tensor",
-      "semantic": "axes",
-      "role": "axes",
-      "dtype": "int32",
-      "narrowing": "checked",
-      "required": false
-    },
-    "squeezed": { "kind": "tensor", "semantic": "squeezed", "role": "output" }
-  },
-  "tunables": { "WORKGROUP_SIZE": 256 },
   "derive": { "axesContract": "not present.axes or ranks.axes == 1" },
-  "constants": { "usesF16": "dtypes.T == \"f16\"" },
   "variants": [
     {
       "id": "copy_vec4",
       "priority": 20,
-      "when": ["axesContract", "numel(shapes.data) > 0", "numel(shapes.data) % 4 == 0", "numel(shapes.data) == numel(shapes.squeezed)", "f16Ok(dtypes.T)"],
-      "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "Squeeze.copyVec4",
-          "source": { "shader": "datamove-flat-copy.wgsl.jinja", "inputs": { "count": "numel(shapes.squeezed) / 4" } },
           "bindings": [
-            {
-              "name": "src",
-              "arg": "data",
-              "semantic": "data",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$vectorScalar"
-            },
-            {
-              "name": "dst",
-              "arg": "squeezed",
-              "semantic": "squeezed",
-              "buffer": { "type": "storage" },
-              "elementType": "$vectorScalar"
-            }
           ],
-          "dispatch": { "gridStride": "numel(shapes.squeezed) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" },
           "viewAlias": [{ "input": "src", "output": "dst" }]
         }
       ]
     },
     {
       "id": "copy",
-      "when": ["axesContract", "numel(shapes.data) == numel(shapes.squeezed)", "f16Ok(dtypes.T)"],
-      "constants": { "scalar": "dtypes.T" },
       "passes": [
         {
           "id": "main",
           "name": "Squeeze",
           "shader": "datamove-elementwise-copy.wgsl.jinja",
           "bindings": [
-            {
-              "name": "x",
-              "arg": "data",
-              "semantic": "data",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$scalar"
-            },
-            {
-              "name": "y",
-              "arg": "squeezed",
-              "semantic": "squeezed",
-              "buffer": { "type": "storage" },
-              "elementType": "$scalar"
-            },
-            {
-              "name": "params",
-              "semantic": "kernel.params",
-              "buffer": { "type": "uniform" },
-              "struct": {
-                "name": "Params",
-                "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.squeezed)" }]
-              }
-            }
           ],
-          "dispatch": { "gridStride": "numel(shapes.squeezed)", "workgroupSize": "tunables.WORKGROUP_SIZE" },
           "viewAlias": [{ "input": "x", "output": "y" }]
         }
       ]

   "domain": "ai.onnx",
   "name": "Squeeze",
   "sinceVersion": 13,
+  "inputs": {
+    "data": { "dtype": "T" },
+    "axes": { "dtype": "S", "rank": 1, "optional": true, "storage": "int32", "narrowing": "checked" }
+  },
+  "outputs": { "squeezed": { "dtype": "T" } },
   "typeConstraints": {
     "T": ["float32", "float16", "int32", "uint32", "int16", "uint8", "int8", "bool"],
     "S": ["int64"]
   },
+  "tunables": { "WORKGROUP_SIZE": { "default": 256 } },
   "derive": { "axesContract": "not present.axes or ranks.axes == 1" },
+  "when": ["axesContract", "numel(shapes.data) == numel(shapes.squeezed)", "f16Ok(dtypes.T)"],
   "variants": [
     {
       "id": "copy_vec4",
       "priority": 20,
+      "when": ["numel(shapes.data) > 0", "numel(shapes.data) % 4 == 0"],
+      "derive": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "Squeeze.copyVec4",
+          "shader": "datamove-flat-copy.wgsl.jinja",
+          "derive": { "count": "numel(shapes.squeezed) / 4" },
           "bindings": [
+            { "arg": "data", "name": "src", "elementType": "$vectorScalar" },
+            { "arg": "squeezed", "name": "dst", "elementType": "$vectorScalar" }
           ],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.squeezed) / 4), (tunables.WORKGROUP_SIZE)), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
+            "y": 1,
+            "z": 1
+          },
           "viewAlias": [{ "input": "src", "output": "dst" }]
         }
       ]
     },
     {
       "id": "copy",
+      "derive": { "scalar": "dtypes.T" },
       "passes": [
         {
           "id": "main",
           "name": "Squeeze",
           "shader": "datamove-elementwise-copy.wgsl.jinja",
           "bindings": [
+            { "arg": "data", "name": "x", "elementType": "$scalar" },
+            { "arg": "squeezed", "name": "y", "elementType": "$scalar" },
+            { "name": "params", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.squeezed)" }] }
           ],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.squeezed)), (tunables.WORKGROUP_SIZE)), min(device.limits.maxComputeWorkgroupsPerDimension, 65535))",
+            "y": 1,
+            "z": 1
+          },
           "viewAlias": [{ "input": "x", "output": "y" }]
         }
       ]

build/webgpu/metadata.json CHANGED Viewed

@@ -1,19 +1,22 @@
 {
   "name": "ai.onnx.Squeeze",
-  "id": "_ai_onnx_squeeze_webgpu_832197a",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "nv9qMEydsiA2N2U6kTkSg/0bUKGXz3olKJQONAe+ALw=",
-      "datamove-elementwise-copy.wgsl.jinja": "J5yC2bAddPiP+odXLgVGS3TJ9jeNsfRTvedKrj/fhZg=",
-      "datamove-flat-copy.wgsl.jinja": "rEpeYMvIZQuQnB2f3lt51jrrklKICsqUTzCJPRTknNk=",
-      "manifest.json": "X1FYKHl8iSCmL4jG2aIzk6pzQGL3QrYp3Jca594J2vQ=",
-      "test.json": "McQKVkPXhYGWdGWzD6vp5gR1x5Cwujt0CkQGow0TGkg="
     }
   },
-  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Squeeze" }
 }

 {
   "name": "ai.onnx.Squeeze",
+  "id": "_ai_onnx_squeeze_webgpu_d1535eb",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "1WC6P+yyucqPFwrdWD/aka7rpPv3lqBN6HF9diF8KhQ=",
+      "datamove-elementwise-copy.wgsl.jinja": "Q1WuZCqcDf9b6cbT4rLilpeWj6327t/bSFmKr6IVBoQ=",
+      "datamove-flat-copy.wgsl.jinja": "g9d62mer2bmHfbScqX5CIT5Zq/Pdinac0W2dH0CD+5s=",
+      "manifest.json": "ZGJbhNJTmqaJQUFcnMmznVL5t66lXFS6aQyFvFXRhT0=",
+      "test.json": "KYRUy5FOrvaucqzu/7scvf/NrMiDWTdWU8e1tyzXwak="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": { "copy_vec4": ["datamove-flat-copy.wgsl.jinja"], "copy": ["datamove-elementwise-copy.wgsl.jinja"] }
+  }
 }

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.Squeeze",
   "fixtureArrays": {
     "bool_axis1_rank4_to_rank3_input_data": [1, 0, 1, 0, 0, 1, 0, 1, 0, 1, 1, 0, 1, 0, 0, 1, 1, 1, 1, 0, 1, 0, 0, 1]
   },
@@ -259,7 +258,7 @@
       "outputs": { "squeezed": { "dtype": "float32", "shape": [3, 4, 5] } },
       "provenance": {
         "source": "cmake/external/onnx/onnx/backend/test/data/node/test_squeeze",
-        "notes": "ONNX int64 metadata/index tensors use framework int32/uint32 slots where representable."
       }
     },
     {
@@ -278,7 +277,7 @@
       "outputs": { "squeezed": { "dtype": "float32", "shape": [1, 3, 5] } },
       "provenance": {
         "source": "cmake/external/onnx/onnx/backend/test/data/node/test_squeeze_negative_axes",
-        "notes": "ONNX int64 metadata/index tensors use framework int32/uint32 slots where representable."
       }
     },
     {
@@ -380,7 +379,7 @@
     {
       "name": "f32_vec4_copy_1024",
       "provenance": {
-        "notes": "Compact vec4-copy sibling for the 1M f32 Squeeze benchmark; keeps the all-unit-axis removal pattern with a small aligned payload."
       },
       "inputs": {
         "data": {

 {
   "fixtureArrays": {
     "bool_axis1_rank4_to_rank3_input_data": [1, 0, 1, 0, 0, 1, 0, 1, 0, 1, 1, 0, 1, 0, 0, 1, 1, 1, 1, 0, 1, 0, 0, 1]
   },
       "outputs": { "squeezed": { "dtype": "float32", "shape": [3, 4, 5] } },
       "provenance": {
         "source": "cmake/external/onnx/onnx/backend/test/data/node/test_squeeze",
+        "notes": "The integer axes tensor exactly represents the source ONNX int64 axis values."
       }
     },
     {
       "outputs": { "squeezed": { "dtype": "float32", "shape": [1, 3, 5] } },
       "provenance": {
         "source": "cmake/external/onnx/onnx/backend/test/data/node/test_squeeze_negative_axes",
+        "notes": "The integer axes tensor exactly represents the source ONNX int64 axis values."
       }
     },
     {
     {
       "name": "f32_vec4_copy_1024",
       "provenance": {
+        "notes": "A small aligned float32 payload exercises vectorized copying while removing all unit axes."
       },
       "inputs": {
         "data": {