sync 91d990483a17

Browse files

Files changed (6) hide show

README.md +11 -7
build/webgpu/bench.json +0 -1
build/webgpu/bitwise-not.wgsl.jinja +16 -16
build/webgpu/manifest.json +31 -49
build/webgpu/metadata.json +10 -7
build/webgpu/test.json +0 -1

README.md CHANGED Viewed

@@ -18,15 +18,15 @@ See the [ONNX `BitwiseNot` spec](https://onnx.ai/onnx/operators/onnx__BitwiseNot
 ## Inputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `X` | `x` | `T` | — | — | Input integer tensor. | required |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `Y` | `y` | `T` | same as `X` | same as `X` | Output tensor with each element bitwise-negated; same shape and dtype as `X`. | required |
 ## Type constraints
@@ -36,7 +36,7 @@ See the [ONNX `BitwiseNot` spec](https://onnx.ai/onnx/operators/onnx__BitwiseNot
 ## 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
@@ -44,10 +44,14 @@ See the [ONNX `BitwiseNot` spec](https://onnx.ai/onnx/operators/onnx__BitwiseNot
 ## 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` | `X` | `T` | — | — | Input integer tensor. | required |
 ## Outputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `y` | `Y` | `T` | same as `x` | same as `x` | Output tensor with each element bitwise-negated; same shape and dtype as `X`. | 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
+```
+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.BitwiseNot",
   "cases": [
     {
       "name": "u32_33m",

 {
   "cases": [
     {
       "name": "u32_33m",

build/webgpu/bitwise-not.wgsl.jinja CHANGED Viewed

@@ -1,30 +1,30 @@
 {% 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;
   }
@@ -34,11 +34,11 @@
 {{ env.wgsl.resourceDeclarations }}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
 {{ flat_index_2d(note="element-limit") }}
   var value = ~x[i];
 {% if logicalDtype == "uint8" %}
-{% if source.vectorized %}
   value = value & vec4<u32>(0xffu);
 {% else %}
   value = value & 0xffu;

 {% 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 dispatch's
+  // per-axis workgroup fold width (outputs > 16.7M elements).
 {% elif note == "limit" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
 {% elif note == "device-axis" %}
+  // The flat dispatch is folded across x/y at a fixed per-axis workgroup
+  // width; 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 dispatch's
+  // per-axis workgroup fold width (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
+  // dispatch's per-axis workgroup fold width.
 {% elif note == "dispatch" %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
 {% endif %}
 {% if bound == "" %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
 {%- elif guardInline %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) { return; }
 {%- else %}
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if ({{ name }} >= {{ bound }}) {
     return;
   }
 {{ env.wgsl.resourceDeclarations }}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
 {{ flat_index_2d(note="element-limit") }}
   var value = ~x[i];
 {% if logicalDtype == "uint8" %}
+{% if vectorizedSpec %}
   value = value & vec4<u32>(0xffu);
 {% else %}
   value = value & 0xffu;

build/webgpu/manifest.json CHANGED Viewed

@@ -2,73 +2,55 @@
   "domain": "ai.onnx",
   "name": "BitwiseNot",
   "sinceVersion": 18,
-  "description": "Applies a bitwise NOT to each element of the input tensor, flipping every bit. The output has the same shape and integer dtype as the input.",
-  "inputs": [{ "role": "X", "dtype": "T", "description": "Input integer tensor." }],
-  "outputs": [
-    {
-      "role": "Y",
-      "dtype": "T",
-      "rank": "ranks.X",
-      "description": "Output tensor with each element bitwise-negated; same shape and dtype as `X`.",
-      "shape": "shapes.X"
-    }
-  ],
   "typeConstraints": { "T": ["uint32", "int32", "int16", "uint8", "int8"] },
-  "args": {
-    "x": { "kind": "tensor", "semantic": "X", "role": "input" },
-    "y": { "kind": "tensor", "semantic": "Y", "role": "output" }
-  },
-  "tunables": { "WORKGROUP_SIZE": 256 },
-  "derive": { "shapeContract": "numel(shapes.X) == numel(shapes.Y)" },
-  "constants": { "logicalDtype": "tensorDtypes.Y" },
-  "bindingSets": {
-    "vector": [
-      { "name": "x", "arg": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$vectorScalar" },
-      { "name": "y", "arg": "y", "buffer": { "type": "storage" }, "elementType": "$vectorScalar" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y) / 4" }] }
-      }
-    ],
-    "scalar": [
-      { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
-      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$T" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.Y)" }] }
-      }
-    ]
-  },
   "variants": [
     {
       "id": "elementwise_vec4",
       "priority": 20,
-      "when": ["numel(shapes.X) % 4 == 0", "shapeContract"],
-      "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "BitwiseNot.vec4",
-          "source": { "shader": "bitwise-not.wgsl.jinja", "inputs": { "vectorized": true } },
-          "bindings": "vector",
-          "dispatch": { "threads": "numel(shapes.Y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "elementwise",
-      "when": "shapeContract",
       "passes": [
         {
           "id": "main",
           "name": "BitwiseNot",
-          "source": { "shader": "bitwise-not.wgsl.jinja", "inputs": { "vectorized": false } },
-          "bindings": "scalar",
-          "dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     }

   "domain": "ai.onnx",
   "name": "BitwiseNot",
   "sinceVersion": 18,
+  "inputs": { "x": { "onnx": "X", "dtype": "T" } },
+  "outputs": { "y": { "onnx": "Y", "dtype": "T", "rank": "ranks.x", "shape": "shapes.x" } },
   "typeConstraints": { "T": ["uint32", "int32", "int16", "uint8", "int8"] },
+  "tunables": { "WORKGROUP_SIZE": { "default": 256 } },
+  "derive": { "shapeContract": "numel(shapes.x) == numel(shapes.y)", "logicalDtype": "tensorDtypes.y" },
+  "when": ["shapeContract"],
   "variants": [
     {
       "id": "elementwise_vec4",
       "priority": 20,
+      "when": ["numel(shapes.x) % 4 == 0"],
+      "derive": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "BitwiseNot.vec4",
+          "shader": "bitwise-not.wgsl.jinja",
+          "derive": { "vectorizedSpec": true },
+          "bindings": [
+            { "arg": "x", "elementType": "$vectorScalar" },
+            { "arg": "y", "elementType": "$vectorScalar" },
+            { "name": "params", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y) / 4" }] }
+          ],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "elementwise",
       "passes": [
         {
           "id": "main",
           "name": "BitwiseNot",
+          "shader": "bitwise-not.wgsl.jinja",
+          "derive": { "vectorizedSpec": false },
+          "bindings": [
+            "x",
+            "y",
+            { "name": "params", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }] }
+          ],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.y)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     }

build/webgpu/metadata.json CHANGED Viewed

@@ -1,18 +1,21 @@
 {
   "name": "ai.onnx.BitwiseNot",
-  "id": "_ai_onnx_bitwisenot_webgpu_32a96bb",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "CczI58x8pA6aakhXdpmw14wRRq8CEArt333TMzvUE5I=",
-      "bitwise-not.wgsl.jinja": "f8h+mARfRrpuDY+hBRyHRhegkFnDEct4lQBZx36dcos=",
-      "manifest.json": "0dDdOkc503LzGY8f0BAs48R9B/tEFd/aiqfgfYE/5/Y=",
-      "test.json": "XY7peYVfnFEA8/3VjLC1K8VgcdVQMfTvctow0rwlmTM="
     }
   },
-  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.BitwiseNot" }
 }

 {
   "name": "ai.onnx.BitwiseNot",
+  "id": "_ai_onnx_bitwisenot_webgpu_853bcee",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "MD1dhXerFPrKL1pwLulPwQoIcIJYxQRi6DZEqLmoETg=",
+      "bitwise-not.wgsl.jinja": "ZtA3TBUEhXt39AUn36XrrbII2ndOJO6W+D3gVBiSMp8=",
+      "manifest.json": "ifbhqCLw5vH+vD1dd5GtcdijLT+JDK1N98R/WnGEiq8=",
+      "test.json": "o049sQp+WvVyGnQxzbNjh0At7j0Yn/3g6v73ROxa1iY="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": { "elementwise_vec4": ["bitwise-not.wgsl.jinja"], "elementwise": ["bitwise-not.wgsl.jinja"] }
+  }
 }

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.BitwiseNot",
   "cases": [
     {
       "name": "int32_signed_ort",

 {
   "cases": [
     {
       "name": "int32_signed_ort",