sync 6fdf6301e2bb

Files changed (5) hide show

README.md CHANGED Viewed

@@ -48,13 +48,13 @@ Attributes and default values (overridable per request):
 - [`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
 - [`bitcast.wgsl.jinja`](build/webgpu/bitcast.wgsl.jinja)
 ## 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.
@@ -68,7 +68,7 @@ Replace each `*Data` placeholder with a typed array containing the corresponding
 import { getKernel } from "@huggingface/kernels";
 const kernel = await getKernel("webgpu-kernels/ai.onnx.BitCast", { version: 1 });
-const { output } = await kernel({ input: { data: inputData, shape: [] } }, {
-  attrs: { to: 6 },
 });
 ```

 - [`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 cases
 - [`bitcast.wgsl.jinja`](build/webgpu/bitcast.wgsl.jinja)
 ## Use with `@huggingface/kernels`
 ```sh
+npm install --save-exact @huggingface/kernels@0.0.1-preview.3
 ```
 Required output shapes and logical data types are inferred from the supplied inputs and attributes; result tensors are allocated automatically.
 import { getKernel } from "@huggingface/kernels";
 const kernel = await getKernel("webgpu-kernels/ai.onnx.BitCast", { version: 1 });
+const { output } = await kernel({ input: { data: inputData, shape: [3] } }, {
+  attrs: { to: 1 },
 });
 ```

build/webgpu/bitcast.wgsl.jinja CHANGED Viewed

@@ -1,3 +1,11 @@
 {{ env.wgsl.resourceDeclarations }}
 {% set vectorized = vectorizedSpec if vectorizedSpec is defined else false %}
@@ -6,12 +14,7 @@
 // their explicit low-byte masking and sign extension.
 @compute @workgroup_size({{ workgroupSizeSpec }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
-  // The flat dispatch is folded across x/y at the device limit, so gid.y
-  // carries the high portion of the slot index.
-  let i = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ workgroupSizeSpec }}u;
-  if (i >= params.count) {
-    return;
-  }
 {% if inScalar == outScalar %}
   output[i] = input[i];
 {% elif inputIsInt8 and outputIsUint8 %}

+{% macro flat_index_2d(workgroupSize, name="i", bound="params.count", guardInline=false) %}
+{% set wgTerm = workgroupSize ~ "u" if workgroupSize is number else workgroupSize %}
+  // 2D-folded flat index: gid.y carries the high bits past the dispatch's
+  // per-axis workgroup fold width.
+  let {{ name }} = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wgTerm }};
+  if ({{ name }} >= {{ bound }}) {
+    return;
+  }{% endmacro %}
 {{ env.wgsl.resourceDeclarations }}
 {% set vectorized = vectorizedSpec if vectorizedSpec is defined else false %}
 // their explicit low-byte masking and sign extension.
 @compute @workgroup_size({{ workgroupSizeSpec }})
 fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+{{ flat_index_2d(workgroupSizeSpec) }}
 {% if inScalar == outScalar %}
   output[i] = input[i];
 {% elif inputIsInt8 and outputIsUint8 %}

build/webgpu/manifest.json CHANGED Viewed

@@ -35,7 +35,7 @@
       "passes": [
         {
           "id": "main",
-          "name": "BitCast.vec4",
           "shader": "bitcast.wgsl.jinja",
           "derive": { "vectorizedSpec": true, "workgroupSizeSpec": "bitcastWorkgroupSize" },
           "bindings": [

       "passes": [
         {
           "id": "main",
+          "name": "BitCast.Vec4",
           "shader": "bitcast.wgsl.jinja",
           "derive": { "vectorizedSpec": true, "workgroupSizeSpec": "bitcastWorkgroupSize" },
           "bindings": [

build/webgpu/metadata.json CHANGED Viewed

@@ -1,6 +1,6 @@
 {
   "name": "ai.onnx.BitCast",
-  "id": "_ai_onnx_bitcast_webgpu_ab905c6",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
@@ -8,14 +8,14 @@
     "algorithm": "sha256",
     "files": {
       "bench.json": "7aNcJ8LB0eyjvFh7jXkwAw1p2vgZejuWw6H9iwqbgqY=",
-      "bitcast.wgsl.jinja": "1vfOUaYUxSx36bWOnKDPyXUPGprDxhPVEv10804xfVI=",
-      "manifest.json": "UplDa7horbJakWMkW0Gkocun6e1kcewHPZAvi6weNHs=",
-      "test.json": "FXueeEbCTZUCVYOfr0NjJrQAoVhhvtwqADXKZI6829c="
     }
   },
-  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
   "webgpu": {
-    "manifestSpec": "2.0",
     "variants": { "slot32_vec4": ["bitcast.wgsl.jinja"], "slot32": ["bitcast.wgsl.jinja"] }
   }
 }

 {
   "name": "ai.onnx.BitCast",
+  "id": "_ai_onnx_bitcast_webgpu_31d09ea",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
     "algorithm": "sha256",
     "files": {
       "bench.json": "7aNcJ8LB0eyjvFh7jXkwAw1p2vgZejuWw6H9iwqbgqY=",
+      "bitcast.wgsl.jinja": "ibrqjbQlHMWRA/FhpRHG1KZup5FFZtRSRIUG9hTcXiQ=",
+      "manifest.json": "h6XBC5d/eqbX9wJxoyVa7PiJ7aNXwEsZCXr91+F9F0c=",
+      "test.json": "TQvH4fE/V77N+0mk4mwR5PxcXsRctMnGBvn1S0LYlqQ="
     }
   },
+  "provenance": { "kernel": { "sha": "6fdf6301e2bbcc2f03bf1eaf493b7ad55ef33afc", "dirty": false } },
   "webgpu": {
+    "manifestSpec": "2.1",
     "variants": { "slot32_vec4": ["bitcast.wgsl.jinja"], "slot32": ["bitcast.wgsl.jinja"] }
   }
 }

build/webgpu/test.json CHANGED Viewed

@@ -314,7 +314,7 @@
     {
       "name": "float32_to_int32_rank7",
       "provenance": {
-        "source": "ONNX BitCast-26 contract and ONNX Runtime CPUExecutionProvider",
         "notes": "BitCast preserves arbitrary tensor rank and only reinterprets same-width element bits."
       },
       "attrs": { "to": 6 },

     {
       "name": "float32_to_int32_rank7",
       "provenance": {
+        "source": "ONNX BitCast-26 contract and ONNX Runtime's CPU provider",
         "notes": "BitCast preserves arbitrary tensor rank and only reinterprets same-width element bits."
       },
       "attrs": { "to": 6 },