sync 91d990483a17

Browse files

Files changed (6) hide show

README.md +14 -10
build/webgpu/bench.json +0 -1
build/webgpu/elementwise-bias-gelu.wgsl.jinja +9 -38
build/webgpu/manifest.json +67 -121
build/webgpu/metadata.json +17 -7
build/webgpu/test.json +4 -5

README.md CHANGED Viewed

@@ -18,16 +18,16 @@ See the [ONNX Runtime `FastGelu` contrib-operator spec](https://github.com/micro
 ## Inputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
-| --- | --- | --- | --- | --- | --- | --- |
-| `X` | `X` | `T` | — | — | Values transformed by FastGelu after adding the optional `bias`. | required |
-| `bias` | `bias` | `T` | `1` | — | Optional 1-D bias added to `X` along the last dimension before the GELU activation. | optional |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
-| --- | --- | --- | --- | --- | --- | --- |
-| `Y` | `Y` | `T` | same as `X` | same as `X` | Output tensor after applying the GELU activation; same shape as `X`. | required |
 ## Type constraints
@@ -37,7 +37,7 @@ See the [ONNX Runtime `FastGelu` contrib-operator spec](https://github.com/micro
 ## 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 `FastGelu` contrib-operator spec](https://github.com/micro
 ## 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` | — | — | Values transformed by FastGelu after adding the optional `bias`. | required |
+| `bias` | `T` | `1` | — | Optional 1-D bias added to `X` along the last dimension before the GELU activation. | optional |
 ## Outputs
+| Name | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- |
+| `Y` | `T` | same as `X` | same as `X` | Output tensor after applying the GELU activation; same shape 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": "com.microsoft.FastGelu",
   "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
   "cases": [
     {

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

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

@@ -1,58 +1,29 @@
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 {% set wg = workgroupSize if workgroupSize is defined else tunables.WORKGROUP_SIZE %}
-// Bias plus GELU, with a specialization-selected tanh or erf approximation.
-// The optional bias is a rank-1 vector broadcast over the innermost (hidden)
-// axis: bias index = element_index % HIDDEN. The `vec4` path requires
-// HIDDEN % 4 == 0 and numel % 4 == 0 so a vec4 group never crosses the hidden
-// axis (the bias slice is then contiguous). `vec4Tail` keeps scalar bindings but
-// evaluates four guarded lanes per invocation, so odd hidden sizes retain the
-// same parallel efficiency without crossing row/bias boundaries. Gelu math and overflow guards match
-// the vectorized unary implementation: tanh saturates to +/-1 by |x|~9, and the
-// erf path uses the same rational approximation as Gelu.
-{% if approximate == "erf" %}
-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;
-  }
-  let sign = select(-1.0, 1.0, x >= 0.0);
-  let t = 1.0 / (1.0 + 0.3275911 * ax);
-  let y = 1.0 - (((((1.061405429 * t - 1.453152027) * t) + 1.421413741) * t - 0.284496736) * t + 0.254829592) * t * exp(-(ax * ax));
-  return sign * y;
-}
-{% else %}
 fn tanh_safe(x: f32) -> f32 {
   if (x > 10.0) { return 1.0; }
   if (x < -10.0) { return -1.0; }
   return tanh(x);
 }
-{% endif %}
 fn gelu_value(v: f32) -> f32 {
-{% if approximate == "erf" %}
-  return 0.5 * v * (1.0 + erf_approx(v * 0.7071067811865476));
-{% else %}
   return 0.5 * v * (1.0 + tanh_safe(0.7978845608028654 * (v + 0.044715 * v * v * v)));
-{% endif %}
 }
 {% if hasBias %}
-const HIDDEN: u32 = {{ hidden }}u;
 {% endif %}
 @compute @workgroup_size({{ wg }})
-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 past the
-  // maxComputeWorkgroupsPerDimension dispatch limit (outputs > 16.7M elements).
-  let i = gid.x + gid.y * nwg.x * {{ wg }}u;
   if (i >= params.count) {
     return;
   }

 {{ env.wgsl.resourceDeclarations }}
 {% set wg = workgroupSize if workgroupSize is defined else tunables.WORKGROUP_SIZE %}
+// Computes GELU, optionally after adding a rank-1 bias broadcast over the
+// innermost axis. A present bias uses element_index % hidden_extent. The `vec4`
+// path requires the innermost extent and element count to be divisible by four;
+// `vec4Tail` instead uses scalar bindings with four guarded lanes.
+// GELU uses the tanh approximation below, with its input clamped in the tails.
 fn tanh_safe(x: f32) -> f32 {
   if (x > 10.0) { return 1.0; }
   if (x < -10.0) { return -1.0; }
   return tanh(x);
 }
 fn gelu_value(v: f32) -> f32 {
   return 0.5 * v * (1.0 + tanh_safe(0.7978845608028654 * (v + 0.044715 * v * v * v)));
 }
 {% if hasBias %}
+const HIDDEN: u32 = {{ hidden | default(0) }}u;
 {% endif %}
 @compute @workgroup_size({{ wg }})
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
   // 2D-folded flat index: gid.y carries the high bits past the
+  // per-axis dispatch fold width (outputs > 16.7M elements).
+  let i = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ wg }}u;
   if (i >= params.count) {
     return;
   }

build/webgpu/manifest.json CHANGED Viewed

@@ -2,120 +2,42 @@
   "domain": "com.microsoft",
   "name": "FastGelu",
   "sinceVersion": 1,
-  "description": "Applies the GELU (Gaussian Error Linear Unit) activation using a tanh approximation: `Y = 0.5 * X * (1 + tanh(0.797885 * X + 0.035677 * X^3))`. An optional `bias` is added to `X` before the activation is computed. This WebGPU package implements float16 and float32; the schema-allowed double and bfloat16 types are not supported.",
-  "inputs": [
-    { "role": "X", "dtype": "T", "description": "Values transformed by FastGelu after adding the optional `bias`." },
-    {
-      "role": "bias",
-      "dtype": "T",
-      "rank": 1,
-      "optional": true,
-      "description": "Optional 1-D bias added to `X` along the last dimension before the GELU activation."
-    }
-  ],
-  "outputs": [
-    {
-      "role": "Y",
-      "dtype": "T",
-      "rank": "ranks.X",
-      "shape": "shapes.X",
-      "description": "Output tensor after applying the GELU activation; same shape as `X`."
-    }
-  ],
   "typeConstraints": { "T": ["float32", "float16"] },
-  "args": {
-    "X": { "kind": "tensor", "semantic": "X", "role": "input" },
-    "bias": { "kind": "tensor", "semantic": "bias", "role": "input", "required": false },
-    "Y": { "kind": "tensor", "semantic": "Y", "role": "output" }
-  },
-  "tunables": { "WORKGROUP_SIZE": 256 },
   "derive": {
     "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
     "storageBufferLimit": "min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
     "workgroupSize": "min(tunables.WORKGROUP_SIZE, deviceWorkgroupCap)",
     "buffersFit": "numel(shapes.X) * dtypeBytes(dtypes.T) <= storageBufferLimit and numel(shapes.Y) * dtypeBytes(dtypes.T) <= storageBufferLimit and (not present.bias or numel(shapes.bias) * dtypeBytes(dtypes.T) <= storageBufferLimit)",
-    "dispatchFits": "numel(shapes.X) <= device.limits.maxComputeWorkgroupsPerDimension * device.limits.maxComputeWorkgroupsPerDimension * workgroupSize * 4",
     "baseOk": "workgroupSize > 0 and ranks.X >= 1 and numel(shapes.X) == numel(shapes.Y) and f16Ok(dtypes.T) and buffersFit and dispatchFits",
     "biasOk": "present.bias and ranks.bias == 1 and dim(shapes.bias, 0) == dim(shapes.X, ranks.X - 1)",
     "noBiasOk": "not present.bias",
-    "vec4Ok": "numel(shapes.X) > 0 and numel(shapes.X) % 4 == 0 and dim(shapes.X, ranks.X - 1) % 4 == 0"
   },
-  "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"", "approximate": "\"tanh\"" },
-  "bindingSets": {
-    "vec4Bias": [
-      {
-        "name": "x",
-        "arg": "X",
-        "semantic": "X",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$vectorScalar"
-      },
-      {
-        "name": "bias",
-        "arg": "bias",
-        "semantic": "bias",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$scalar",
-        "length": "$hidden"
-      },
-      { "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" }] }
-      }
-    ],
-    "vec4": [
-      {
-        "name": "x",
-        "arg": "X",
-        "semantic": "X",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$vectorScalar"
-      },
-      { "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" }] }
-      }
-    ],
-    "scalarBias": [
-      { "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.X)" }] }
-      }
-    ],
-    "scalar": [
-      { "name": "x", "arg": "X", "semantic": "X", "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)" }] }
-      }
-    ]
   },
   "variants": [
     {
       "id": "vec4_bias",
       "priority": 30,
-      "when": ["baseOk", "biasOk", "vec4Ok"],
-      "constants": {
         "vec4": true,
         "vec4Tail": false,
         "hasBias": true,
@@ -127,31 +49,39 @@
           "id": "main",
           "name": "FastGelu.vec4Bias",
           "shader": "elementwise-bias-gelu.wgsl.jinja",
-          "bindings": "vec4Bias",
-          "dispatch": { "threads": "numel(shapes.X) / 4", "workgroupSize": "workgroupSize" }
         }
       ]
     },
     {
       "id": "vec4_no_bias",
       "priority": 25,
-      "when": ["baseOk", "noBiasOk", "vec4Ok"],
-      "constants": { "vec4": true, "vec4Tail": false, "hasBias": false, "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "FastGelu.vec4",
           "shader": "elementwise-bias-gelu.wgsl.jinja",
-          "bindings": "vec4",
-          "dispatch": { "threads": "numel(shapes.X) / 4", "workgroupSize": "workgroupSize" }
         }
       ]
     },
     {
       "id": "vec4_tail_bias",
       "priority": 20,
-      "when": ["baseOk", "biasOk", "numel(shapes.X) > 0"],
-      "constants": {
         "vec4": false,
         "vec4Tail": true,
         "hasBias": true,
@@ -162,31 +92,39 @@
           "id": "main",
           "name": "FastGelu.vec4TailBias",
           "shader": "elementwise-bias-gelu.wgsl.jinja",
-          "bindings": "scalarBias",
-          "dispatch": { "threads": "ceilDiv(numel(shapes.X), 4)", "workgroupSize": "workgroupSize" }
         }
       ]
     },
     {
       "id": "vec4_tail_no_bias",
       "priority": 15,
-      "when": ["baseOk", "noBiasOk", "numel(shapes.X) > 0"],
-      "constants": { "vec4": false, "vec4Tail": true, "hasBias": false },
       "passes": [
         {
           "id": "main",
           "name": "FastGelu.vec4Tail",
           "shader": "elementwise-bias-gelu.wgsl.jinja",
-          "bindings": "scalar",
-          "dispatch": { "threads": "ceilDiv(numel(shapes.X), 4)", "workgroupSize": "workgroupSize" }
         }
       ]
     },
     {
       "id": "scalar_bias",
       "priority": 10,
-      "when": ["baseOk", "biasOk", "true"],
-      "constants": {
         "vec4": false,
         "vec4Tail": false,
         "hasBias": true,
@@ -197,23 +135,31 @@
           "id": "main",
           "name": "FastGelu.scalarBias",
           "shader": "elementwise-bias-gelu.wgsl.jinja",
-          "bindings": "scalarBias",
-          "dispatch": { "threads": "numel(shapes.X)", "workgroupSize": "workgroupSize" }
         }
       ]
     },
     {
       "id": "scalar_no_bias",
       "priority": 0,
-      "when": ["baseOk", "noBiasOk", "true"],
-      "constants": { "vec4": false, "vec4Tail": false, "hasBias": false },
       "passes": [
         {
           "id": "main",
           "name": "FastGelu.scalar",
           "shader": "elementwise-bias-gelu.wgsl.jinja",
-          "bindings": "scalar",
-          "dispatch": { "threads": "numel(shapes.X)", "workgroupSize": "workgroupSize" }
         }
       ]
     }

   "domain": "com.microsoft",
   "name": "FastGelu",
   "sinceVersion": 1,
+  "inputs": { "X": { "dtype": "T" }, "bias": { "dtype": "T", "rank": 1, "optional": true } },
+  "outputs": { "Y": { "dtype": "T", "rank": "ranks.X", "shape": "shapes.X" } },
   "typeConstraints": { "T": ["float32", "float16"] },
+  "tunables": { "WORKGROUP_SIZE": { "default": 256 } },
   "derive": {
     "deviceWorkgroupCap": "min(device.limits.maxComputeInvocationsPerWorkgroup, device.limits.maxComputeWorkgroupSizeX)",
     "storageBufferLimit": "min(device.limits.maxStorageBufferBindingSize, device.limits.maxBufferSize)",
     "workgroupSize": "min(tunables.WORKGROUP_SIZE, deviceWorkgroupCap)",
     "buffersFit": "numel(shapes.X) * dtypeBytes(dtypes.T) <= storageBufferLimit and numel(shapes.Y) * dtypeBytes(dtypes.T) <= storageBufferLimit and (not present.bias or numel(shapes.bias) * dtypeBytes(dtypes.T) <= storageBufferLimit)",
+    "dispatchFits": "numel(shapes.X) <= min(device.limits.maxComputeWorkgroupsPerDimension, 65535) * min(device.limits.maxComputeWorkgroupsPerDimension, 65535) * workgroupSize * 4",
     "baseOk": "workgroupSize > 0 and ranks.X >= 1 and numel(shapes.X) == numel(shapes.Y) and f16Ok(dtypes.T) and buffersFit and dispatchFits",
     "biasOk": "present.bias and ranks.bias == 1 and dim(shapes.bias, 0) == dim(shapes.X, ranks.X - 1)",
     "noBiasOk": "not present.bias",
+    "vec4Ok": "numel(shapes.X) > 0 and numel(shapes.X) % 4 == 0 and dim(shapes.X, ranks.X - 1) % 4 == 0",
+    "scalar": "dtypes.T",
+    "approximate": "\"tanh\""
   },
+  "when": ["baseOk"],
+  "bindings": {
+    "x": { "arg": "X", "buffer": "read-only-storage", "elementType": "$vectorScalar" },
+    "y": { "arg": "Y", "buffer": "storage", "elementType": "$vectorScalar" },
+    "params": { "buffer": "uniform", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.X) / 4" }] },
+    "x_2": { "arg": "X", "name": "x", "buffer": "read-only-storage", "elementType": "$scalar" },
+    "y_2": { "arg": "Y", "name": "y", "buffer": "storage", "elementType": "$scalar" },
+    "params_2": {
+      "name": "params",
+      "buffer": "uniform",
+      "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.X)" }]
+    }
   },
   "variants": [
     {
       "id": "vec4_bias",
       "priority": 30,
+      "when": ["biasOk", "vec4Ok"],
+      "derive": {
         "vec4": true,
         "vec4Tail": false,
         "hasBias": true,
           "id": "main",
           "name": "FastGelu.vec4Bias",
           "shader": "elementwise-bias-gelu.wgsl.jinja",
+          "bindings": ["x", { "arg": "bias", "elementType": "$scalar", "length": "$hidden" }, "y", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.X) / 4), (workgroupSize)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.X) / 4), (workgroupSize)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "vec4_no_bias",
       "priority": 25,
+      "when": ["noBiasOk", "vec4Ok"],
+      "derive": { "vec4": true, "vec4Tail": false, "hasBias": false, "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "FastGelu.vec4",
           "shader": "elementwise-bias-gelu.wgsl.jinja",
+          "bindings": ["x", "y", "params"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.X) / 4), (workgroupSize)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.X) / 4), (workgroupSize)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "vec4_tail_bias",
       "priority": 20,
+      "when": ["biasOk", "numel(shapes.X) > 0"],
+      "derive": {
         "vec4": false,
         "vec4Tail": true,
         "hasBias": true,
           "id": "main",
           "name": "FastGelu.vec4TailBias",
           "shader": "elementwise-bias-gelu.wgsl.jinja",
+          "bindings": ["x_2", "bias", "y_2", "params_2"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.X), 4)), (workgroupSize)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.X), 4)), (workgroupSize)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "vec4_tail_no_bias",
       "priority": 15,
+      "when": ["noBiasOk", "numel(shapes.X) > 0"],
+      "derive": { "vec4": false, "vec4Tail": true, "hasBias": false },
       "passes": [
         {
           "id": "main",
           "name": "FastGelu.vec4Tail",
           "shader": "elementwise-bias-gelu.wgsl.jinja",
+          "bindings": ["x_2", "y_2", "params_2"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.X), 4)), (workgroupSize)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.X), 4)), (workgroupSize)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "scalar_bias",
       "priority": 10,
+      "when": ["biasOk", "true"],
+      "derive": {
         "vec4": false,
         "vec4Tail": false,
         "hasBias": true,
           "id": "main",
           "name": "FastGelu.scalarBias",
           "shader": "elementwise-bias-gelu.wgsl.jinja",
+          "bindings": ["x_2", "bias", "y_2", "params_2"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.X)), (workgroupSize)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.X)), (workgroupSize)), 65535)",
+            "z": 1
+          }
         }
       ]
     },
     {
       "id": "scalar_no_bias",
       "priority": 0,
+      "when": ["noBiasOk", "true"],
+      "derive": { "vec4": false, "vec4Tail": false, "hasBias": false },
       "passes": [
         {
           "id": "main",
           "name": "FastGelu.scalar",
           "shader": "elementwise-bias-gelu.wgsl.jinja",
+          "bindings": ["x_2", "y_2", "params_2"],
+          "dispatch": {
+            "x": "min(ceilDiv((numel(shapes.X)), (workgroupSize)), 65535)",
+            "y": "ceilDiv(ceilDiv((numel(shapes.X)), (workgroupSize)), 65535)",
+            "z": 1
+          }
         }
       ]
     }

build/webgpu/metadata.json CHANGED Viewed

@@ -1,18 +1,28 @@
 {
   "name": "com.microsoft.FastGelu",
-  "id": "_com_microsoft_fastgelu_webgpu_efaff66",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "m93bKUUMLGZxG/F5G6AFYMyDIKSE1Pk4hOyYNKrO+y0=",
-      "elementwise-bias-gelu.wgsl.jinja": "Eg8N2jJCMce+IsYNcCzuxvs8CSv7yTQXdorvv+c0L58=",
-      "manifest.json": "WckeD7rVTEJ1Lb6FDdpYC9n//UUq4+QWwd7vyAGc5m8=",
-      "test.json": "shIb34pVGz4K2kmzCn/yOmL6OuRbXD+UaRJ/PADEFtM="
     }
   },
-  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/com.microsoft.FastGelu" }
 }

 {
   "name": "com.microsoft.FastGelu",
+  "id": "_com_microsoft_fastgelu_webgpu_da3aabd",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "uob2jmkWOUzFyVqQy0sJhB7JkQGvh+5ux8/7f+O6DYA=",
+      "elementwise-bias-gelu.wgsl.jinja": "OW2nqbCYsSLMd8w5lKcNx4zjdkHi4i2kB+FFAKyHhFg=",
+      "manifest.json": "nHE7yhM3Q5obzXWntT8sXryEtqYicMVDVSqdP+kwbLs=",
+      "test.json": "oH8XR6itu84Ti3GZbKE8ImRXTl4N2njDh1qRw0+HhCs="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "vec4_bias": ["elementwise-bias-gelu.wgsl.jinja"],
+      "vec4_no_bias": ["elementwise-bias-gelu.wgsl.jinja"],
+      "vec4_tail_bias": ["elementwise-bias-gelu.wgsl.jinja"],
+      "vec4_tail_no_bias": ["elementwise-bias-gelu.wgsl.jinja"],
+      "scalar_bias": ["elementwise-bias-gelu.wgsl.jinja"],
+      "scalar_no_bias": ["elementwise-bias-gelu.wgsl.jinja"]
+    }
+  }
 }

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "com.microsoft.FastGelu",
   "fixtureArrays": {
     "ort_float16_hidden8_with_bias_input_X": [0.8, -0.5, 0, 1, 1.3, 2.1, -0.2, 1.1, 0.5, 0.2, 0.3, -0.6, 3.1, 2.2, -1.1, 0]
   },
@@ -77,7 +76,7 @@
     {
       "name": "f32_scalar_no_bias_zero_sequence",
       "provenance": {
-        "notes": "Bias-free twin of ort_float32_zero_sequence_with_bias. An empty X fails the numel > 0 guard both vec4 routes carry, so the plain scalar no-bias kernel is the only one left."
       },
       "inputs": { "X": { "dtype": "float32", "shape": [1, 0, 4], "data": { "kind": "values", "values": [] } } },
       "outputs": { "Y": { "dtype": "float32", "shape": [1, 0, 4], "data": { "kind": "values", "values": [] } } }
@@ -300,7 +299,7 @@
       "provenance": {
         "source": "onnxruntime/test/contrib_ops/fastgelu_op_test.cc",
         "test": "FastGeluTest.FastGeluWithoutBiasFloat32",
-        "notes": "Scalar-path companion for FastGelu's near-zero x/2 subnormal behavior."
       },
       "inputs": {
         "X": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1e-40, -1e-40, 1e-38] } }
@@ -323,7 +322,7 @@
       "provenance": {
         "source": "onnxruntime/test/contrib_ops/fastgelu_op_test.cc",
         "test": "FastGeluTest.FastGeluWithBiasFloat32",
-        "notes": "Bias-path companion: zero bias reduces FastGelu to the same near-zero x/2 behavior, but exercises the vec4 bias-broadcast kernel variant."
       },
       "inputs": {
         "X": {
@@ -354,7 +353,7 @@
       "provenance": {
         "source": "onnxruntime/test/contrib_ops/fastgelu_op_test.cc",
         "test": "FastGeluTest.FastGeluWithBiasFloat32",
-        "notes": "Scalar bias-broadcast companion for FastGelu subnormal linear-region behavior with zero bias."
       },
       "inputs": {
         "X": { "dtype": "float32", "shape": [1, 3], "data": { "kind": "values", "values": [1e-40, -1e-40, 1e-38] } },

 {
   "fixtureArrays": {
     "ort_float16_hidden8_with_bias_input_X": [0.8, -0.5, 0, 1, 1.3, 2.1, -0.2, 1.1, 0.5, 0.2, 0.3, -0.6, 3.1, 2.2, -1.1, 0]
   },
     {
       "name": "f32_scalar_no_bias_zero_sequence",
       "provenance": {
+        "notes": "An empty bias-free input bypasses both vec4 routes and selects the scalar no-bias kernel."
       },
       "inputs": { "X": { "dtype": "float32", "shape": [1, 0, 4], "data": { "kind": "values", "values": [] } } },
       "outputs": { "Y": { "dtype": "float32", "shape": [1, 0, 4], "data": { "kind": "values", "values": [] } } }
       "provenance": {
         "source": "onnxruntime/test/contrib_ops/fastgelu_op_test.cc",
         "test": "FastGeluTest.FastGeluWithoutBiasFloat32",
+        "notes": "Near-zero inputs exercise FastGelu's subnormal x/2 behavior on the scalar path."
       },
       "inputs": {
         "X": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [1e-40, -1e-40, 1e-38] } }
       "provenance": {
         "source": "onnxruntime/test/contrib_ops/fastgelu_op_test.cc",
         "test": "FastGeluTest.FastGeluWithBiasFloat32",
+        "notes": "Zero bias isolates near-zero x/2 behavior on the vec4 bias-broadcast path."
       },
       "inputs": {
         "X": {
       "provenance": {
         "source": "onnxruntime/test/contrib_ops/fastgelu_op_test.cc",
         "test": "FastGeluTest.FastGeluWithBiasFloat32",
+        "notes": "Zero bias isolates subnormal linear-region behavior on the scalar bias-broadcast path."
       },
       "inputs": {
         "X": { "dtype": "float32", "shape": [1, 3], "data": { "kind": "values", "values": [1e-40, -1e-40, 1e-38] } },