sync 91d990483a17

Browse files

Files changed (6) hide show

README.md +15 -11
build/webgpu/bench.json +0 -1
build/webgpu/bias-split-gelu.wgsl.jinja +23 -28
build/webgpu/manifest.json +49 -128
build/webgpu/metadata.json +14 -7
build/webgpu/test.json +5 -6

README.md CHANGED Viewed

@@ -12,22 +12,22 @@ tags:
 ## Description
-Adds `bias` to `X`, splits the last dimension in half, then multiplies the left half elementwise by the GELU activation of the right half, producing an output with half the hidden dimension. Used as a fused activation in diffusion models.
 See the [ONNX Runtime `BiasSplitGelu` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.BiasSplitGelu) for the reference semantics.
 ## Inputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
-| --- | --- | --- | --- | --- | --- | --- |
-| `X` | `X` | `T` | `3` | — | Input tensor of shape `(N, S, D)`, where `N` is the batch size, `S` is the spatial size, and `D` is the hidden dimension. | required |
-| `bias` | `bias` | `T` | `1` | — | 1-D bias tensor of length D, matching the hidden dimension of X. | required |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
-| --- | --- | --- | --- | --- | --- | --- |
-| `Y` | `Y` | `T` | `3` | derived; see description | Output tensor of shape `(N, S, D/2)` after the split-GELU activation. | required |
 ## Type constraints
@@ -37,7 +37,7 @@ See the [ONNX Runtime `BiasSplitGelu` contrib-operator spec](https://github.com/
 ## 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
@@ -45,10 +45,14 @@ See the [ONNX Runtime `BiasSplitGelu` contrib-operator spec](https://github.com/
 ## 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.

 ## Description
+Adds `bias` to `X`, splits the last dimension in half, then multiplies the left half elementwise by the GELU activation of the right half, producing an output with half the hidden dimension.
 See the [ONNX Runtime `BiasSplitGelu` contrib-operator spec](https://github.com/microsoft/onnxruntime/blob/main/docs/ContribOperators.md#com.microsoft.BiasSplitGelu) for the reference semantics.
 ## Inputs
+| Name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `X` | `T` | `3` | — | Input tensor of shape `(N, S, D_in)`, where `N` is the batch size, `S` is the spatial size, and `D_in` is an even hidden dimension. | required |
+| `bias` | `T` | `1` | — | 1-D bias tensor of length `D_in`, matching the input hidden dimension. | required |
 ## Outputs
+| Name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `Y` | `T` | `3` | derived | Output tensor of shape `(N, S, D_out)`, where `D_out = D_in / 2`. | 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.BiasSplitGelu",
   "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
   "cases": [
     {

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

build/webgpu/bias-split-gelu.wgsl.jinja CHANGED Viewed

@@ -1,56 +1,51 @@
 {% 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.BiasSplitGelu GEGLU gate:
-//   The innermost (hidden) axis of (X + bias) has size 2*D. Split it into
-//   value = [.., 0:D] and gate = [.., D:2D], then
-//   Y[.., d] = value[.., d] * gelu_erf(gate[.., d])
 //   gelu_erf(g) = 0.5 * g * (1 + erf(g / sqrt2)).
-// bias is a rank-1 vector [2*D] broadcast over the hidden axis. The output
-// hidden axis is D (= HIDDEN below); the input hidden axis is 2*D (TWICE).
 // erf_approx uses an overflow-safe rational approximation.
 fn erf_approx(x: f32) -> f32 {
   let ax = abs(x);
   // The polynomial has a small nonzero floor near zero. Use erf(x) ~=
   // 2/sqrt(pi)*x below 2^-20 to preserve erf(0) == 0, odd symmetry, and the
-  // correctly rounded f32 result. The exactly representable threshold keeps
-  // scalar and vector branching identical. NaN falls through to the polynomial
-  // and propagates.
   if (ax < 9.5367431640625e-7) {
     return 1.1283791670955126 * x;
   }
@@ -64,14 +59,14 @@ fn gelu_value(v: f32) -> f32 {
   return 0.5 * v * (1.0 + erf_approx(v * 0.7071067811865476));
 }
-const HIDDEN: u32 = {{ hidden }}u;   // D: output hidden axis
-const TWICE: u32 = {{ twice }}u;     // 2*D: input hidden axis
 {% if vec4Tail %}
 const GROUPS_PER_ROW: u32 = (HIDDEN + 3u) / 4u;
 {% endif %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
-fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups) nwg: vec3<u32>) {
 {{ flat_index_2d() }}
 {% if vec4Tail %}
   let row = i / GROUPS_PER_ROW;

 {% 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.BiasSplitGelu GEGLU gate:
+//   The innermost axis of (X + bias) has size D_in = 2*D_out. Split it into
+//   value = [.., 0:D_out] and gate = [.., D_out:D_in], then
+//   Y[.., d] = value[.., d] * gelu_erf(gate[.., d]).
 //   gelu_erf(g) = 0.5 * g * (1 + erf(g / sqrt2)).
+// `bias` is a rank-1 vector of length D_in broadcast over the hidden axis.
 // erf_approx uses an overflow-safe rational approximation.
 fn erf_approx(x: f32) -> f32 {
   let ax = abs(x);
   // The polynomial has a small nonzero floor near zero. Use erf(x) ~=
   // 2/sqrt(pi)*x below 2^-20 to preserve erf(0) == 0, odd symmetry, and the
+  // small-input linear behavior. The threshold is exactly representable in
+  // f32. NaN falls through to the polynomial and propagates.
   if (ax < 9.5367431640625e-7) {
     return 1.1283791670955126 * x;
   }
   return 0.5 * v * (1.0 + erf_approx(v * 0.7071067811865476));
 }
+const HIDDEN: u32 = {{ hidden }}u;   // D_out
+const TWICE: u32 = {{ twice }}u;     // D_in
 {% if vec4Tail %}
 const GROUPS_PER_ROW: u32 = (HIDDEN + 3u) / 4u;
 {% endif %}
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
 {{ flat_index_2d() }}
 {% if vec4Tail %}
   let row = i / GROUPS_PER_ROW;

build/webgpu/manifest.json CHANGED Viewed

@@ -2,45 +2,40 @@
   "domain": "com.microsoft",
   "name": "BiasSplitGelu",
   "sinceVersion": 1,
-  "description": "Adds `bias` to `X`, splits the last dimension in half, then multiplies the left half elementwise by the GELU activation of the right half, producing an output with half the hidden dimension. Used as a fused activation in diffusion models.",
-  "inputs": [
-    {
-      "role": "X",
-      "dtype": "T",
-      "rank": 3,
-      "description": "Input tensor of shape `(N, S, D)`, where `N` is the batch size, `S` is the spatial size, and `D` is the hidden dimension."
     },
-    {
-      "role": "bias",
-      "dtype": "T",
-      "rank": 1,
-      "description": "1-D bias tensor of length D, matching the hidden dimension of X."
-    }
-  ],
-  "outputs": [
-    {
-      "role": "Y",
-      "dtype": "T",
-      "rank": 3,
-      "shape": "[dim(shapes.X, 0), dim(shapes.X, 1), dim(shapes.X, 2) / 2]",
-      "description": "Output tensor of shape `(N, S, D/2)` after the split-GELU activation."
     }
-  ],
-  "typeConstraints": { "T": ["float32", "float16"] },
-  "args": {
-    "X": { "kind": "tensor", "semantic": "X", "role": "input" },
-    "bias": { "kind": "tensor", "semantic": "bias", "role": "input" },
-    "Y": { "kind": "tensor", "semantic": "Y", "role": "output" }
   },
-  "tunables": { "WORKGROUP_SIZE": 64 },
   "variants": [
     {
       "id": "vec4",
       "priority": 30,
-      "when": ["ranks.X == 3", "ranks.Y == 3", "dim(shapes.Y, 0) == dim(shapes.X, 0)", "dim(shapes.Y, 1) == dim(shapes.X, 1)", "dim(shapes.X, 2) % 2 == 0", "dim(shapes.Y, 2) == dim(shapes.X, 2) / 2", "f16Ok(dtypes.T)", "ranks.bias == 1", "dim(shapes.bias, 0) == dim(shapes.X, 2)", "numel(shapes.Y) > 0", "numel(shapes.Y) % 4 == 0", "dim(shapes.Y, 2) % 4 == 0"],
-      "constants": {
         "scalar": "dtypes.T",
-        "usesF16": "dtypes.T == \"f16\"",
         "vec4": true,
         "vec4Tail": false,
         "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
@@ -53,49 +48,25 @@
           "name": "BiasSplitGelu.vec4",
           "shader": "bias-split-gelu.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": "$twice"
-            },
-            {
-              "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.Y) / 4" }]
-              }
-            }
           ],
-          "dispatch": { "threads": "numel(shapes.Y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     },
     {
       "id": "vec4_tail",
       "priority": 20,
-      "when": ["ranks.X == 3", "ranks.Y == 3", "dim(shapes.Y, 0) == dim(shapes.X, 0)", "dim(shapes.Y, 1) == dim(shapes.X, 1)", "dim(shapes.X, 2) % 2 == 0", "dim(shapes.Y, 2) == dim(shapes.X, 2) / 2", "f16Ok(dtypes.T)", "ranks.bias == 1", "dim(shapes.bias, 0) == dim(shapes.X, 2)", "numel(shapes.Y) > 0"],
-      "constants": {
         "scalar": "dtypes.T",
-        "usesF16": "dtypes.T == \"f16\"",
         "vec4": false,
         "vec4Tail": true,
         "hidden": "dim(shapes.X, ranks.X - 1) / 2 if dim(shapes.X, ranks.X - 1) > 0 else 1",
@@ -106,42 +77,11 @@
           "id": "main",
           "name": "BiasSplitGelu.vec4Tail",
           "shader": "bias-split-gelu.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": "$twice"
-            },
-            { "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.Y) / (dim(shapes.X, ranks.X - 1) / 2)) * ceilDiv(dim(shapes.X, ranks.X - 1) / 2, 4)"
-                  }
-                ]
-              }
-            }
-          ],
           "dispatch": {
-            "threads": "(numel(shapes.Y) / (dim(shapes.X, ranks.X - 1) / 2)) * ceilDiv(dim(shapes.X, ranks.X - 1) / 2, 4)",
-            "workgroupSize": "tunables.WORKGROUP_SIZE"
           }
         }
       ]
@@ -149,10 +89,9 @@
     {
       "id": "scalar",
       "priority": 0,
-      "when": ["ranks.X == 3", "ranks.Y == 3", "dim(shapes.Y, 0) == dim(shapes.X, 0)", "dim(shapes.Y, 1) == dim(shapes.X, 1)", "dim(shapes.X, 2) % 2 == 0", "dim(shapes.Y, 2) == dim(shapes.X, 2) / 2", "f16Ok(dtypes.T)", "ranks.bias == 1", "dim(shapes.bias, 0) == dim(shapes.X, 2)", "true"],
-      "constants": {
         "scalar": "dtypes.T",
-        "usesF16": "dtypes.T == \"f16\"",
         "vec4": false,
         "vec4Tail": false,
         "hidden": "dim(shapes.X, ranks.X - 1) / 2 if dim(shapes.X, ranks.X - 1) > 0 else 1",
@@ -163,30 +102,12 @@
           "id": "main",
           "name": "BiasSplitGelu.scalar",
           "shader": "bias-split-gelu.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"
-            },
-            { "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.Y)" }] }
-            }
-          ],
-          "dispatch": { "threads": "numel(shapes.Y)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     }

   "domain": "com.microsoft",
   "name": "BiasSplitGelu",
   "sinceVersion": 1,
+  "inputs": { "X": { "dtype": "T", "rank": 3 }, "bias": { "dtype": "T", "rank": 1 } },
+  "outputs": { "Y": { "dtype": "T", "rank": 3, "shape": "[dim(shapes.X, 0), dim(shapes.X, 1), dim(shapes.X, 2) / 2]" } },
+  "typeConstraints": { "T": ["float32", "float16"] },
+  "tunables": { "WORKGROUP_SIZE": { "default": 64 } },
+  "when": ["ranks.X == 3", "ranks.Y == 3", "dim(shapes.Y, 0) == dim(shapes.X, 0)", "dim(shapes.Y, 1) == dim(shapes.X, 1)", "dim(shapes.X, 2) % 2 == 0", "dim(shapes.Y, 2) == dim(shapes.X, 2) / 2", "f16Ok(dtypes.T)", "ranks.bias == 1", "dim(shapes.bias, 0) == dim(shapes.X, 2)"],
+  "bindings": {
+    "x": { "arg": "X", "buffer": "read-only-storage", "elementType": "$scalar" },
+    "bias": { "buffer": "read-only-storage", "elementType": "$scalar", "length": "$twice" },
+    "y_2": { "arg": "Y", "name": "y", "buffer": "storage", "elementType": "$scalar" },
+    "params_2": {
+      "name": "params",
+      "buffer": "uniform",
+      "struct": [
+        {
+          "name": "count",
+          "type": "u32",
+          "value": "(numel(shapes.Y) / (dim(shapes.X, ranks.X - 1) / 2)) * ceilDiv(dim(shapes.X, ranks.X - 1) / 2, 4)"
+        }
+      ]
     },
+    "bias_2": { "name": "bias", "buffer": "read-only-storage", "elementType": "$scalar" },
+    "params_3": {
+      "name": "params",
+      "buffer": "uniform",
+      "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.Y)" }]
     }
   },
   "variants": [
     {
       "id": "vec4",
       "priority": 30,
+      "when": ["numel(shapes.Y) > 0", "numel(shapes.Y) % 4 == 0", "dim(shapes.Y, 2) % 4 == 0"],
+      "derive": {
         "scalar": "dtypes.T",
         "vec4": true,
         "vec4Tail": false,
         "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"",
           "name": "BiasSplitGelu.vec4",
           "shader": "bias-split-gelu.wgsl.jinja",
           "bindings": [
+            "x",
+            "bias",
+            { "arg": "Y", "name": "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": "vec4_tail",
       "priority": 20,
+      "when": ["numel(shapes.Y) > 0"],
+      "derive": {
         "scalar": "dtypes.T",
         "vec4": false,
         "vec4Tail": true,
         "hidden": "dim(shapes.X, ranks.X - 1) / 2 if dim(shapes.X, ranks.X - 1) > 0 else 1",
           "id": "main",
           "name": "BiasSplitGelu.vec4Tail",
           "shader": "bias-split-gelu.wgsl.jinja",
+          "bindings": ["x", "bias", "y_2", "params_2"],
           "dispatch": {
+            "x": "min(ceilDiv(((numel(shapes.Y) / (dim(shapes.X, ranks.X - 1) / 2)) * ceilDiv(dim(shapes.X, ranks.X - 1) / 2, 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv(((numel(shapes.Y) / (dim(shapes.X, ranks.X - 1) / 2)) * ceilDiv(dim(shapes.X, ranks.X - 1) / 2, 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
           }
         }
       ]
     {
       "id": "scalar",
       "priority": 0,
+      "when": ["true"],
+      "derive": {
         "scalar": "dtypes.T",
         "vec4": false,
         "vec4Tail": false,
         "hidden": "dim(shapes.X, ranks.X - 1) / 2 if dim(shapes.X, ranks.X - 1) > 0 else 1",
           "id": "main",
           "name": "BiasSplitGelu.scalar",
           "shader": "bias-split-gelu.wgsl.jinja",
+          "bindings": ["x", "bias_2", "y_2", "params_3"],
+          "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,25 @@
 {
   "name": "com.microsoft.BiasSplitGelu",
-  "id": "_com_microsoft_biassplitgelu_webgpu_e56edec",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "oYxkYz1Ukonqiav8OJSz76L0uidTkGg049tpmfnC0Ms=",
-      "bias-split-gelu.wgsl.jinja": "YGnskVbzCTa0y/SN+O6Wq0Yp8SWpUav6Q930/EHEvsw=",
-      "manifest.json": "i40t5ZCHmqMoQrSWSWZlC3X/ERSmyoa5z4PkrzsWP/w=",
-      "test.json": "JVElcOkuNl49j2nfKWBgotONhDTNVT7ZbWTDW+5Of/Q="
     }
   },
-  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/com.microsoft.BiasSplitGelu" }
 }

 {
   "name": "com.microsoft.BiasSplitGelu",
+  "id": "_com_microsoft_biassplitgelu_webgpu_c7117c6",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "9D5pXr6GQoiKWOaIZonasPZV750s7EBoeCc18Fjh6LE=",
+      "bias-split-gelu.wgsl.jinja": "NdmUm+O7ZCH0nHgT9xuUVfsMcX0TKQLoz2edprOP4Eg=",
+      "manifest.json": "RhauSBIjgU0p1ri0W74BHIABvKxAMnSg9MfBB++0c0E=",
+      "test.json": "jeqt2QA6oOrulPiXyZ1nXe7lvFKPg1me9DSnpzLx7lk="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "vec4": ["bias-split-gelu.wgsl.jinja"],
+      "vec4_tail": ["bias-split-gelu.wgsl.jinja"],
+      "scalar": ["bias-split-gelu.wgsl.jinja"]
+    }
+  }
 }

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "com.microsoft.BiasSplitGelu",
   "cases": [
     {
       "name": "dispatch_cliff_vec4_over_16M",
@@ -47,7 +46,7 @@
       "provenance": {
         "source": "onnxruntime/test/contrib_ops/bias_split_gelu_op_test.cc",
         "test": "BiasSplitGeluTest.BiasSplitGeluTest_HiddenSize_5120",
-        "notes": "Compact deterministic projection of ORT's batch=2, sequence_length=1 shape coverage."
       },
       "inputs": {
         "X": {
@@ -79,7 +78,7 @@
       "provenance": {
         "source": "onnxruntime/test/contrib_ops/bias_split_gelu_op_test.cc",
         "test": "BiasSplitGeluTest.BiasSplitGeluTest_HiddenSize_10240",
-        "notes": "Compact deterministic projection of ORT's batch=1, sequence_length=2 shape coverage."
       },
       "inputs": {
         "X": {
@@ -111,7 +110,7 @@
       "provenance": {
         "source": "onnxruntime/test/contrib_ops/bias_split_gelu_op_test.cc",
         "test": "BiasSplitGeluTest.BiasSplitGeluTest_HiddenSize_6144",
-        "notes": "Compact deterministic projection of ORT's batch=2, sequence_length=3 shape coverage with output hidden divisible by 4."
       },
       "inputs": {
         "X": {
@@ -228,7 +227,7 @@
     {
       "name": "scalar_path_dispatch_fold_over_16M",
       "provenance": {
-        "notes": "D=6 (twice=12, 6%4!=0) forces the scalar variant; numel(Y)=16776966 > 65535*256 forces the scalar-branch 2D dispatch fold (gid.y high bits), an untested path (dispatch_cliff_vec4_over_16M covers only the vec4 branch). Validates the i=gid.x+gid.y*nwg.x*WG reconstruction and i>=count tail guard on the last rows."
       },
       "inputs": {
         "X": {
@@ -243,7 +242,7 @@
     {
       "name": "scalar_path_odd_hidden_d6",
       "provenance": {
-        "notes": "twice=12, D=6 (6%4!=0) routes to the scalar variant; small case verifying scalar split-index math (value half [0:6], gate half [6:12]) at a hidden size not covered by existing D=3 scalar cases. Reference supplies truth."
       },
       "inputs": {
         "X": {

 {
   "cases": [
     {
       "name": "dispatch_cliff_vec4_over_16M",
       "provenance": {
         "source": "onnxruntime/test/contrib_ops/bias_split_gelu_op_test.cc",
         "test": "BiasSplitGeluTest.BiasSplitGeluTest_HiddenSize_5120",
+        "notes": "A compact deterministic ORT case with batch 2 and sequence length 1."
       },
       "inputs": {
         "X": {
       "provenance": {
         "source": "onnxruntime/test/contrib_ops/bias_split_gelu_op_test.cc",
         "test": "BiasSplitGeluTest.BiasSplitGeluTest_HiddenSize_10240",
+        "notes": "A compact deterministic ORT case with batch 1 and sequence length 2."
       },
       "inputs": {
         "X": {
       "provenance": {
         "source": "onnxruntime/test/contrib_ops/bias_split_gelu_op_test.cc",
         "test": "BiasSplitGeluTest.BiasSplitGeluTest_HiddenSize_6144",
+        "notes": "A compact deterministic ORT case with batch 2, sequence length 3, and output width divisible by four."
       },
       "inputs": {
         "X": {
     {
       "name": "scalar_path_dispatch_fold_over_16M",
       "provenance": {
+        "notes": "Output width 6 is not divisible by four, selecting the scalar variant. With 16,776,966 output elements, the two-dimensional dispatch fold must reconstruct the linear index from both grid axes and guard the final partial row."
       },
       "inputs": {
         "X": {
     {
       "name": "scalar_path_odd_hidden_d6",
       "provenance": {
+        "notes": "With input width 12 and output width 6, the unaligned output uses scalar indexing. Values [0:6] form the value half and [6:12] form the gate half."
       },
       "inputs": {
         "X": {