sync 91d990483a17

Browse files

Files changed (6) hide show

README.md +15 -11
build/webgpu/bench.json +0 -1
build/webgpu/bias-add.wgsl.jinja +22 -33
build/webgpu/manifest.json +34 -115
build/webgpu/metadata.json +10 -7
build/webgpu/test.json +0 -1

README.md CHANGED Viewed

@@ -18,17 +18,17 @@ See the [ONNX Runtime `BiasAdd` contrib-operator spec](https://github.com/micros
 ## Inputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
-| --- | --- | --- | --- | --- | --- | --- |
-| `X` | `X` | `T` | `3` | — | Input tensor of shape `(N, S, C)`: batch size `N`, spatial size `S`, and `C` channels. | required |
-| `bias` | `bias` | `T` | `1` | — | 1-D bias vector of length C, broadcast-added along the channel dimension. | required |
-| `skip` | `skip` | `T` | `3` | — | Residual tensor with the same `(N, S, C)` shape as `X`, added after the bias. | required |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
-| --- | --- | --- | --- | --- | --- | --- |
-| `Y` | `Y` | `T` | `3` | same as `X` | Output tensor of shape `(N, S, C)`: the elementwise sum `X + bias + skip`. | required |
 ## Type constraints
@@ -38,7 +38,7 @@ See the [ONNX Runtime `BiasAdd` contrib-operator spec](https://github.com/micros
 ## 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
@@ -46,10 +46,14 @@ See the [ONNX Runtime `BiasAdd` contrib-operator spec](https://github.com/micros
 ## 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 | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `X` | `T` | `3` | — | Input tensor of shape `(N, S, C)`: batch size `N`, spatial size `S`, and `C` channels. | required |
+| `bias` | `T` | `1` | — | 1-D bias vector of length C, broadcast-added along the channel dimension. | required |
+| `skip` | `T` | `3` | — | Residual tensor with the same `(N, S, C)` shape as `X`, added after the bias. | required |
 ## Outputs
+| Name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `Y` | `T` | `3` | same as `X` | Output tensor of shape `(N, S, C)`: the elementwise sum `X + bias + skip`. | 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": "com.microsoft.BiasAdd",
   "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
   "cases": [
     {

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

build/webgpu/bias-add.wgsl.jinja CHANGED Viewed

@@ -1,64 +1,53 @@
 {% 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;
-{% if source.itemsPerInvocation is defined %}
   // Tail-safe scalar x4 keeps vector-like dispatch density without requiring
   // the logical tensor length (or its storage binding) to be vec4 aligned.
-  let begin = invocation * {{ source.itemsPerInvocation }}u;
-  let end = min(begin + {{ source.itemsPerInvocation }}u, params.count);
   for (var i = begin; i < end; i = i + 1u) {
-{%- else %}
-  let i = invocation;
-  if (i >= params.count) {
-    return;
-  }
-{%- endif %}
-{% endmacro %}
 {% macro flat_tail_close() %}
-{% if source.itemsPerInvocation is defined %}
   }
-{% endif %}
 {% endmacro %}
 {% 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;
   }
 {%- endif %}
 {% endmacro %}
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 // com.microsoft.BiasAdd : Y = X + bias + skip
@@ -75,7 +64,7 @@ const HIDDEN: u32 = {{ hidden }}u;
 {% if vec4 %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
 {{ flat_index_2d() }}
   let base = i * 4u;
   let xv = vec4<f32>(f32(x[base]), f32(x[base + 1u]), f32(x[base + 2u]), f32(x[base + 3u]));

 {% 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 %}
+  let begin = invocation * {{ itemsPerInvocation }}u;
+  let end = min(begin + {{ itemsPerInvocation }}u, params.count);
   for (var i = begin; i < end; i = i + 1u) {
+{%- endmacro %}
 {% macro flat_tail_close() %}
   }
 {% endmacro %}
 {% 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;
   }
 {%- endif %}
 {% endmacro %}
 {{ env.wgsl.resourceDeclarations }}
 // com.microsoft.BiasAdd : Y = X + bias + skip
 {% if vec4 %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
 {{ flat_index_2d() }}
   let base = i * 4u;
   let xv = vec4<f32>(f32(x[base]), f32(x[base + 1u]), f32(x[base + 2u]), f32(x[base + 3u]));

build/webgpu/manifest.json CHANGED Viewed

@@ -2,146 +2,65 @@
   "domain": "com.microsoft",
   "name": "BiasAdd",
   "sinceVersion": 1,
-  "description": "Adds a 1-D `bias` (broadcast over the channel dimension) to input `X`, then adds the residual tensor `skip` elementwise. All three tensors share the same channel count `C`; `X` and `skip` have shape `(N, S, C)`.",
-  "inputs": [
-    {
-      "role": "X",
-      "dtype": "T",
-      "rank": 3,
-      "description": "Input tensor of shape `(N, S, C)`: batch size `N`, spatial size `S`, and `C` channels."
-    },
-    {
-      "role": "bias",
-      "dtype": "T",
-      "rank": 1,
-      "description": "1-D bias vector of length C, broadcast-added along the channel dimension."
-    },
-    {
-      "role": "skip",
-      "dtype": "T",
-      "rank": 3,
-      "description": "Residual tensor with the same `(N, S, C)` shape as `X`, added after the bias."
-    }
-  ],
-  "outputs": [
-    {
-      "role": "Y",
-      "dtype": "T",
-      "rank": 3,
-      "shape": "shapes.X",
-      "description": "Output tensor of shape `(N, S, C)`: the elementwise sum `X + bias + skip`."
-    }
-  ],
-  "typeConstraints": { "T": ["float32", "float16"] },
-  "args": {
-    "X": { "kind": "tensor", "semantic": "X", "role": "input" },
-    "bias": { "kind": "tensor", "semantic": "bias", "role": "input" },
-    "skip": { "kind": "tensor", "semantic": "skip", "role": "input" },
-    "Y": { "kind": "tensor", "semantic": "Y", "role": "output" }
-  },
-  "tunables": { "WORKGROUP_SIZE": 256 },
-  "constants": {
-    "scalar": "dtypes.T",
-    "usesF16": "dtypes.T == \"f16\"",
-    "hidden": "dim(shapes.X, ranks.X - 1) if dim(shapes.X, ranks.X - 1) > 0 else 1"
   },
   "variants": [
     {
       "id": "vec4",
       "priority": 30,
-      "when": ["ranks.X == 3", "ranks.skip == 3", "ranks.Y == 3", "sameShape(shapes.X, shapes.skip)", "sameShape(shapes.X, shapes.Y)", "f16Ok(dtypes.T)", "ranks.bias == 1", "dim(shapes.bias, 0) == dim(shapes.X, 2)", "numel(shapes.X) > 0", "numel(shapes.X) % 4 == 0", "dim(shapes.X, 2) % 4 == 0"],
-      "constants": { "vec4": true, "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "BiasAdd.vec4",
           "shader": "bias-add.wgsl.jinja",
           "bindings": [
-            {
-              "name": "x",
-              "arg": "X",
-              "semantic": "X",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$scalar"
-            },
-            {
-              "name": "bias",
-              "arg": "bias",
-              "semantic": "bias",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$scalar",
-              "length": "$hidden"
-            },
-            {
-              "name": "skip",
-              "arg": "skip",
-              "semantic": "skip",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$scalar"
-            },
-            {
-              "name": "y",
-              "arg": "Y",
-              "semantic": "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.X) / 4" }]
-              }
-            }
           ],
-          "dispatch": { "threads": "numel(shapes.X) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "scalar",
       "priority": 0,
-      "when": ["ranks.X == 3", "ranks.skip == 3", "ranks.Y == 3", "sameShape(shapes.X, shapes.skip)", "sameShape(shapes.X, shapes.Y)", "f16Ok(dtypes.T)", "ranks.bias == 1", "dim(shapes.bias, 0) == dim(shapes.X, 2)"],
-      "constants": { "vec4": false },
       "passes": [
         {
           "id": "main",
           "name": "BiasAdd.scalar",
           "shader": "bias-add.wgsl.jinja",
-          "source": { "inputs": { "itemsPerInvocation": 4 } },
           "bindings": [
-            {
-              "name": "x",
-              "arg": "X",
-              "semantic": "X",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$scalar"
-            },
-            {
-              "name": "bias",
-              "arg": "bias",
-              "semantic": "bias",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$scalar"
-            },
-            {
-              "name": "skip",
-              "arg": "skip",
-              "semantic": "skip",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$scalar"
-            },
-            { "name": "y", "arg": "Y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$scalar" },
-            {
-              "name": "params",
-              "semantic": "kernel.params",
-              "buffer": { "type": "uniform" },
-              "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.X)" }] }
-            }
           ],
-          "dispatch": { "threads": "ceilDiv(numel(shapes.X), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     }

   "domain": "com.microsoft",
   "name": "BiasAdd",
   "sinceVersion": 1,
+  "inputs": {
+    "X": { "dtype": "T", "rank": 3 },
+    "bias": { "dtype": "T", "rank": 1 },
+    "skip": { "dtype": "T", "rank": 3 }
   },
+  "outputs": { "Y": { "dtype": "T", "rank": 3, "shape": "shapes.X" } },
+  "typeConstraints": { "T": ["float32", "float16"] },
+  "tunables": { "WORKGROUP_SIZE": { "default": 256 } },
+  "derive": { "scalar": "dtypes.T", "hidden": "dim(shapes.X, ranks.X - 1) if dim(shapes.X, ranks.X - 1) > 0 else 1" },
+  "when": ["ranks.X == 3", "ranks.skip == 3", "ranks.Y == 3", "sameShape(shapes.X, shapes.skip)", "sameShape(shapes.X, shapes.Y)", "f16Ok(dtypes.T)", "ranks.bias == 1", "dim(shapes.bias, 0) == dim(shapes.X, 2)"],
+  "bindings": { "x": { "arg": "X", "buffer": "read-only-storage", "elementType": "$scalar" } },
   "variants": [
     {
       "id": "vec4",
       "priority": 30,
+      "when": ["numel(shapes.X) > 0", "numel(shapes.X) % 4 == 0", "dim(shapes.X, 2) % 4 == 0"],
+      "derive": { "vec4": true, "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "BiasAdd.vec4",
           "shader": "bias-add.wgsl.jinja",
           "bindings": [
+            "x",
+            { "arg": "bias", "elementType": "$scalar", "length": "$hidden" },
+            "skip",
+            { "arg": "Y", "name": "y", "elementType": "$vectorScalar" },
+            { "name": "params", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.X) / 4" }] }
           ],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.X) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.X) / 4), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "scalar",
       "priority": 0,
+      "derive": { "vec4": false },
       "passes": [
         {
           "id": "main",
           "name": "BiasAdd.scalar",
           "shader": "bias-add.wgsl.jinja",
+          "derive": { "itemsPerInvocation": 4 },
           "bindings": [
+            "x",
+            "bias",
+            "skip",
+            { "arg": "Y", "name": "y", "elementType": "$scalar" },
+            { "name": "params", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.X)" }] }
           ],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.X), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.X), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     }

build/webgpu/metadata.json CHANGED Viewed

@@ -1,18 +1,21 @@
 {
   "name": "com.microsoft.BiasAdd",
-  "id": "_com_microsoft_biasadd_webgpu_40b1ba1",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "W2cQSIIm7hI8MbHD6Cro9DGB/ByB035ELwNigWwGHYw=",
-      "bias-add.wgsl.jinja": "bZPMYyEnjQ+7kuft3Vjj/ZfXZq2hV5QcMREEyyiGt8w=",
-      "manifest.json": "0KqWQgIXxxbypIfz2fQy91LPsxiEoov4j8V52QlixsE=",
-      "test.json": "PLZQbRu0fhmrRSakHOzPjTiZL0LYBLFxmBdw7gd31m4="
     }
   },
-  "provenance": { "kernel": { "sha": "c928d21e6cc1310861cba3bafb75f5f679ecf5f3", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/com.microsoft.BiasAdd" }
 }

 {
   "name": "com.microsoft.BiasAdd",
+  "id": "_com_microsoft_biasadd_webgpu_7687b51",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "RdL1BQwfXYch113DNTUHHlTuPCDk1Pp6wKa2sV9SeY4=",
+      "bias-add.wgsl.jinja": "w8g/0fT1IBapNvGjVIocSxND3JpDsMvxJQYZHQJOAXw=",
+      "manifest.json": "iwu1wkLNtggvjZU/fzw2dq+3bcm+q/KfLIYAUpManhY=",
+      "test.json": "yECzeltqwOXVNMeShcml5Bz3B2Amx3fmBX47DQCPG70="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": { "vec4": ["bias-add.wgsl.jinja"], "scalar": ["bias-add.wgsl.jinja"] }
+  }
 }

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "com.microsoft.BiasAdd",
   "cases": [
     {
       "name": "dispatch_cliff_scalar_f32",

 {
   "cases": [
     {
       "name": "dispatch_cliff_scalar_f32",