sync 91d990483a17

Browse files

Files changed (7) hide show

README.md +11 -7
build/webgpu/bench.json +0 -1
build/webgpu/manifest.json +92 -48
build/webgpu/metadata.json +15 -8
build/webgpu/test.json +3 -4
build/webgpu/unary-scalar.wgsl.jinja +29 -21
build/webgpu/unary-vec4.wgsl.jinja +37 -20

README.md CHANGED Viewed

@@ -18,15 +18,15 @@ See the [ONNX `Mish` spec](https://onnx.ai/onnx/operators/onnx__Mish.html) for t
 ## Inputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `X` | `x` | `T` | — | — | Values transformed elementwise by the Mish activation. | required |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `Y` | `y` | `T` | same as `X` | same as `X` | Output tensor with the same shape as `X`, with Mish applied elementwise. | required |
 ## Type constraints
@@ -36,7 +36,7 @@ See the [ONNX `Mish` spec](https://onnx.ai/onnx/operators/onnx__Mish.html) for t
 ## 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 `Mish` spec](https://onnx.ai/onnx/operators/onnx__Mish.html) for t
 ## 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` | — | — | Values transformed elementwise by the Mish activation. | required |
 ## Outputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `y` | `Y` | `T` | same as `x` | same as `x` | Output tensor with the same shape as `X`, with Mish applied elementwise. | 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.Mish",
   "cases": [
     {
       "name": "1m_f32",

 {
   "cases": [
     {
       "name": "1m_f32",

build/webgpu/manifest.json CHANGED Viewed

@@ -2,28 +2,16 @@
   "domain": "ai.onnx",
   "name": "Mish",
   "sinceVersion": 18,
-  "description": "Applies the Mish activation function elementwise: `mish(x) = x * tanh(ln(1 + e^x))`. Mish is a self-regularized non-monotonic activation that smoothly gates its input using a softplus-based tanh.",
-  "inputs": [{ "role": "X", "dtype": "T", "description": "Values transformed elementwise by the Mish activation." }],
-  "outputs": [
-    {
-      "role": "Y",
-      "dtype": "T",
-      "rank": "ranks.X",
-      "description": "Output tensor with the same shape as `X`, with Mish applied elementwise.",
-      "shape": "shapes.X"
-    }
-  ],
   "typeConstraints": { "T": ["float32", "float16"] },
-  "args": {
-    "x": { "kind": "tensor", "semantic": "X", "role": "input" },
-    "y": { "kind": "tensor", "semantic": "Y", "role": "output" }
-  },
-  "tunables": { "WORKGROUP_SIZE": 256 },
   "variants": [
     {
       "id": "same_layout_vec4",
       "when": ["numel(shapes.x) > 0", "numel(shapes.x) % 4 == 0", "numel(shapes.x) == numel(shapes.y)", "f16Ok(dtypes.T)"],
-      "constants": {
         "scalar": "dtypes.T",
         "usesF16": "dtypes.T == \"f16\"",
         "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\""
@@ -32,47 +20,103 @@
         {
           "id": "main",
           "name": "Mish.vec4",
-          "source": { "shader": "unary-vec4.wgsl.jinja", "inputs": { "op": "\"mish\"" } },
-          "bindings": [
-            { "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" }]
-              }
-            }
-          ],
-          "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
-      ],
-      "priority": 20
     },
     {
       "id": "elementwise",
       "when": ["numel(shapes.x) == numel(shapes.y)", "f16Ok(dtypes.T)"],
-      "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "Mish",
-          "source": { "shader": "unary-scalar.wgsl.jinja", "inputs": { "op": "\"mish\"", "itemsPerInvocation": 4 } },
-          "bindings": [
-            { "name": "x", "arg": "x", "buffer": { "type": "read-only-storage" }, "elementType": "$scalar" },
-            { "name": "y", "arg": "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": "ceilDiv(numel(shapes.y), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     }
-  ]
 }

   "domain": "ai.onnx",
   "name": "Mish",
   "sinceVersion": 18,
+  "inputs": { "x": { "onnx": "X", "dtype": "T" } },
+  "outputs": { "y": { "onnx": "Y", "dtype": "T", "rank": "ranks.x", "shape": "shapes.x" } },
   "typeConstraints": { "T": ["float32", "float16"] },
+  "tunables": { "WORKGROUP_SIZE": { "default": 256 } },
   "variants": [
     {
       "id": "same_layout_vec4",
+      "priority": 20,
       "when": ["numel(shapes.x) > 0", "numel(shapes.x) % 4 == 0", "numel(shapes.x) == numel(shapes.y)", "f16Ok(dtypes.T)"],
+      "derive": {
         "scalar": "dtypes.T",
         "usesF16": "dtypes.T == \"f16\"",
         "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\""
         {
           "id": "main",
           "name": "Mish.vec4",
+          "shader": "unary-vec4.wgsl.jinja",
+          "derive": {
+            "op": "\"mish\"",
+            "vec4PerThread": "4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1"
+          },
+          "bindings": ["x", "y", "params_unary"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.y) / 4, 4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.y) / 4, 4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
+      ]
     },
     {
       "id": "elementwise",
       "when": ["numel(shapes.x) == numel(shapes.y)", "f16Ok(dtypes.T)"],
+      "derive": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
       "passes": [
         {
           "id": "main",
           "name": "Mish",
+          "shader": "unary-scalar.wgsl.jinja",
+          "derive": { "op": "\"mish\"", "itemsPerInvocation": 4 },
+          "bindings": ["x_2", "y_2", "params_2_unary"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(numel(shapes.y), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(numel(shapes.y), 4)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
+        }
+      ]
+    },
+    {
+      "id": "same_layout_vec4_tail",
+      "priority": 19,
+      "when": ["numel(shapes.x) > 4", "numel(shapes.x) % 4 != 0", "numel(shapes.x) == numel(shapes.y)", "f16Ok(dtypes.T)"],
+      "derive": {
+        "scalar": "dtypes.T",
+        "usesF16": "dtypes.T == \"f16\"",
+        "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\""
+      },
+      "passes": [
+        {
+          "id": "bulk",
+          "name": "Mish.vec4Bulk",
+          "shader": "unary-vec4.wgsl.jinja",
+          "derive": {
+            "op": "\"mish\"",
+            "vec4PerThread": "4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1"
+          },
+          "bindings": ["x", "y", "params_unary_tail"],
+          "dispatch": {
+            "x": "min(ceilDiv((ceilDiv(floor(numel(shapes.y) / 4), 4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((ceilDiv(floor(numel(shapes.y) / 4), 4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1)), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
+        },
+        {
+          "id": "tail",
+          "name": "Mish.tail",
+          "shader": "unary-scalar.wgsl.jinja",
+          "derive": { "op": "\"mish\"", "itemsPerInvocation": 4, "tailOnly": true },
+          "bindings": ["x_2", "y_2", "params_2_unary_tail"],
+          "dispatch": {
+            "x": "min(ceilDiv((1), (tunables.WORKGROUP_SIZE)), 65535)",
+            "y": "ceilDiv(ceilDiv((1), (tunables.WORKGROUP_SIZE)), 65535)",
+            "z": 1
+          }
         }
       ]
     }
+  ],
+  "bindings": {
+    "x": { "buffer": "read-only-storage", "elementType": "$vectorScalar" },
+    "y": { "buffer": "storage", "elementType": "$vectorScalar" },
+    "params_unary": {
+      "buffer": "uniform",
+      "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y) / 4" }],
+      "name": "params"
+    },
+    "x_2": { "buffer": "read-only-storage", "name": "x", "elementType": "$scalar" },
+    "y_2": { "buffer": "storage", "name": "y", "elementType": "$scalar" },
+    "params_2_unary": {
+      "buffer": "uniform",
+      "name": "params",
+      "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }]
+    },
+    "params_unary_tail": {
+      "buffer": "uniform",
+      "struct": [{ "name": "count", "type": "u32", "value": "floor(numel(shapes.y) / 4)" }],
+      "name": "params"
+    },
+    "params_2_unary_tail": {
+      "buffer": "uniform",
+      "name": "params",
+      "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }]
+    }
+  }
 }

build/webgpu/metadata.json CHANGED Viewed

@@ -1,19 +1,26 @@
 {
   "name": "ai.onnx.Mish",
-  "id": "_ai_onnx_mish_webgpu_a99258a",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "faDpALYIY3uK5S1wPVMgs/oHhjLDA9SG3vsgQ+Ft7Ds=",
-      "manifest.json": "gjrNDsfxKYT9/Za4zImkT0t/ZzxPQ13mFYJOvZuRB7I=",
-      "test.json": "jWKrZonQPkjiI5K5te3ZdyIZIxACG7fMVoAvKQjVg2Y=",
-      "unary-scalar.wgsl.jinja": "j5Dt18TaEi25ZySk2z4iEyQcRtkAEYQlOoO60mqG668=",
-      "unary-vec4.wgsl.jinja": "UOWANhQMot9qLN3iIhH89q5NrRkwCD7lAPQ5Fx1bE9Y="
     }
   },
-  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Mish" }
 }

 {
   "name": "ai.onnx.Mish",
+  "id": "_ai_onnx_mish_webgpu_096f2a8",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "TfLj3+DuyE9y+n5bGsOjK1xu3f0uDeBcIAhyvYzR5rU=",
+      "manifest.json": "XKRUFo4lxoSn+9aekUCsUGHOcbdj9tv9Hg+rYA2lXvY=",
+      "test.json": "OUDrD3ZrcYkRW2c77hqlGo/d6Ws8hsUgNk/hMv1pz3c=",
+      "unary-scalar.wgsl.jinja": "z8hXeS1n3GN14/DcKymOFL+ZIqEcu5GJeOjnyypRWJU=",
+      "unary-vec4.wgsl.jinja": "e1wb44e3IKm1LIkdMYM87pnMkTOzqU32SvJPVN7EG64="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "same_layout_vec4": ["unary-vec4.wgsl.jinja"],
+      "elementwise": ["unary-scalar.wgsl.jinja"],
+      "same_layout_vec4_tail": ["unary-scalar.wgsl.jinja", "unary-vec4.wgsl.jinja"]
+    }
+  }
 }

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.Mish",
   "cases": [
     {
       "name": "f32_values",
@@ -51,7 +50,7 @@
       "provenance": {
         "source": "cmake/external/onnx/onnx/backend/test/data/node/test_mish",
         "test": "test_mish",
-        "notes": "Scalar-path companion to the vec4 negative-tail stability case: non-multiple-of-4 shape must still use log1p-style softplus for tiny finite outputs."
       },
       "inputs": {
         "x": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-19.0, -18.0, -17.0] } }
@@ -83,7 +82,7 @@
       "provenance": {
         "source": "cmake/external/onnx/onnx/backend/test/data/node/test_mish",
         "test": "test_mish",
-        "notes": "Scalar-path companion for Mish values whose softplus term is subnormal but whose final product is finite."
       },
       "inputs": {
         "x": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-80.0, -88.0, -90.0] } }
@@ -120,7 +119,7 @@
     {
       "name": "f32_nonfinite_endpoints",
       "provenance": {
-        "notes": "ORT CPU validated nonfinite edge: -Infinity follows the literal formula into NaN, +Infinity stays +Infinity."
       },
       "inputs": {
         "x": {

 {
   "cases": [
     {
       "name": "f32_values",
       "provenance": {
         "source": "cmake/external/onnx/onnx/backend/test/data/node/test_mish",
         "test": "test_mish",
+        "notes": "A non-multiple-of-four shape exercises log1p-style softplus on the scalar path for tiny finite negative-tail outputs."
       },
       "inputs": {
         "x": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-19.0, -18.0, -17.0] } }
       "provenance": {
         "source": "cmake/external/onnx/onnx/backend/test/data/node/test_mish",
         "test": "test_mish",
+        "notes": "On the scalar path, Mish inputs with a subnormal softplus term must still produce the finite expected product."
       },
       "inputs": {
         "x": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-80.0, -88.0, -90.0] } }
     {
       "name": "f32_nonfinite_endpoints",
       "provenance": {
+        "notes": "Nonfinite edge: -Infinity follows the literal formula into NaN, +Infinity stays +Infinity."
       },
       "inputs": {
         "x": {

build/webgpu/unary-scalar.wgsl.jinja CHANGED Viewed

@@ -1,50 +1,58 @@
 {% 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;
   // 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) {
 {%- endmacro %}
 {% macro flat_tail_close() %}
   }
 {% endmacro %}
-// Scalar unary fallback. Each branch retains the operation's numeric hardening,
-// including Payne-Hanek trigonometric range reduction and NaN/overflow guards.
 {% if usesF16 %}
 enable f16;
 {% endif %}
 {{ env.wgsl.resourceDeclarations }}
 fn tanh_safe(x: f32) -> f32 {
-  // The hardware tanh evaluates (e^2x - 1)/(e^2x + 1), which is Inf/Inf = NaN
-  // once e^2x overflows (here softplus(v) ~ v grows without bound, and overflows
-  // to +Inf for v > ~88). tanh saturates to 1 within f32 precision by x ~ 9, so
-  // clamp the tail; this keeps mish(v) -> v for large v instead of NaN.
   if (x > 10.0) { return 1.0; }
   if (x < -10.0) { return -1.0; }
-  // The hardware tanh also collapses to 0 (or a single ULP step) for tiny x,
-  // because its internal e^2x rounds back to 1.0; tanh(x) == x to full f32
-  // precision well before |x| ~ 1e-4 (the x^3/3 term is < 1 ULP), so return x
-  // directly. Without this the small softplus tail of mish flushes to zero.
   if (x > -1.0e-4 && x < 1.0e-4) { return x; }
   return tanh(x);
-}
 fn log1p_small_positive(y: f32) -> f32 {
-  // log1p(y) for small positive y via series, avoiding the catastrophic
-  // cancellation of log(1.0 + y) where 1.0 + y rounds back to 1.0 (mirrors the
-  // Softplus kernel's helper).
   return y * (1.0 + y * (-0.5 + y * (0.3333333333333333 + y * (-0.25 + y * 0.2))));
 }
 fn mish_softplus(v: f32) -> f32 {
   // Numerically stable softplus = log(1 + e^v). The naive log(1.0 + exp(v))
   // collapses to 0 once exp(v) underflows below the 1.0 + x rounding step
-  // (v <~ -16), which zeros mish's small finite negative tail. Match the CPU
-  // log1p form: series for the tiny tail, identity-shifted log otherwise.
   if (v > 20.0) { return v; }
   if (v < -8.0) { return log1p_small_positive(exp(v)); }
   return log(1.0 + exp(-abs(v))) + max(v, 0.0);

 {% 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 %}
+{% if tailOnly is defined and tailOnly %}
+  // Tail of a vec4 bulk pass: lane zero alone covers the elements the packed
+  // pass left behind, from the last multiple of the item count to the end.
+  if (invocation != 0u) {
+    return;
+  }
+  let begin = params.count - params.count % {{ itemsPerInvocation }}u;
+  let end = params.count;
+{% else %}
+  let begin = invocation * {{ itemsPerInvocation }}u;
+  let end = min(begin + {{ itemsPerInvocation }}u, params.count);
+{% endif %}
   for (var i = begin; i < end; i = i + 1u) {
 {%- endmacro %}
 {% macro flat_tail_close() %}
   }
 {% endmacro %}
+// Scalar unary elementwise implementation. Specialization emits only the
+// selected operation and any numerical helper it requires.
 {% if usesF16 %}
 enable f16;
 {% endif %}
 {{ env.wgsl.resourceDeclarations }}
+{% macro emit_tanh_safe() %}
 fn tanh_safe(x: f32) -> f32 {
+  // tanh rounds to its saturated value for these tails in f32. Return that
+  // value directly, including for infinite input, before invoking the builtin.
   if (x > 10.0) { return 1.0; }
   if (x < -10.0) { return -1.0; }
+  // For tiny |x|, return x directly to preserve its sign and magnitude without
+  // relying on backend-specific builtin behavior near zero.
   if (x > -1.0e-4 && x < 1.0e-4) { return x; }
   return tanh(x);
+}{% endmacro %}
+{{ emit_tanh_safe() }}
 fn log1p_small_positive(y: f32) -> f32 {
+  // log1p(y) for small positive y via series, avoiding cancellation in
+  // log(1.0 + y) when 1.0 + y rounds back to 1.0.
   return y * (1.0 + y * (-0.5 + y * (0.3333333333333333 + y * (-0.25 + y * 0.2))));
 }
 fn mish_softplus(v: f32) -> f32 {
   // Numerically stable softplus = log(1 + e^v). The naive log(1.0 + exp(v))
   // collapses to 0 once exp(v) underflows below the 1.0 + x rounding step
+  // (v <~ -16), which zeros Mish's small finite negative tail. Use a series for
+  // the tiny tail and the identity-shifted logarithm otherwise.
   if (v > 20.0) { return v; }
   if (v < -8.0) { return log1p_small_positive(exp(v)); }
   return log(1.0 + exp(-abs(v))) + max(v, 0.0);

build/webgpu/unary-vec4.wgsl.jinja CHANGED Viewed

@@ -1,34 +1,29 @@
-// Loads and stores vec4<T> (128 bits) while retaining scalar per-component
-// arithmetic, including per-component helper calls for guard-heavy operations.
-{% if usesF16 %}
-enable f16;
-{% endif %}
 {{ env.wgsl.resourceDeclarations }}
 fn tanh_safe(x: f32) -> f32 {
-  // The hardware tanh evaluates (e^2x - 1)/(e^2x + 1), which is Inf/Inf = NaN
-  // once e^2x overflows (x > ~44, and at x == +Inf). tanh saturates to +/-1
-  // within f32 precision by |x| ~ 9, so clamp the tail to the exact saturated
-  // value.
   if (x > 10.0) { return 1.0; }
   if (x < -10.0) { return -1.0; }
-  // The hardware tanh also collapses to 0 (or a single ULP step) for tiny x,
-  // because its internal e^2x rounds back to 1.0; tanh(x) == x to full f32
-  // precision well before |x| ~ 1e-4 (the x^3/3 term is < 1 ULP), so return x
-  // directly. Without this the small softplus tail of mish flushes to zero.
   if (x > -1.0e-4 && x < 1.0e-4) { return x; }
   return tanh(x);
-}
 fn mish_log1p_small_positive(y: f32) -> f32 {
   // log1p(y) via series for small positive y, avoiding the log(1.0 + y)
-  // cancellation where 1.0 + y rounds back to 1.0 (mirrors the Softplus helper).
   return y * (1.0 + y * (-0.5 + y * (0.3333333333333333 + y * (-0.25 + y * 0.2))));
 }
 fn mish_softplus(v: f32) -> f32 {
   // Numerically stable softplus = log(1 + e^v). The naive log(1.0 + exp(v))
   // collapses to 0 once exp(v) underflows below the 1.0 + x rounding step
-  // (v <~ -16), which zeros mish's small finite negative tail. Match the CPU
-  // log1p form: series for the tiny tail, identity-shifted log otherwise.
   if (v > 20.0) { return v; }
   if (v < -8.0) { return mish_log1p_small_positive(exp(v)); }
   return log(1.0 + exp(-abs(v))) + max(v, 0.0);
@@ -38,15 +33,37 @@ fn mish_value(v: f32) -> f32 {
   return v * tanh_safe(softplus);
 }
 @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 i = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
   if (i >= params.count) {
     return;
   }
   let xv = x[i];
   let fv = vec4<f32>(xv);
   y[i] = {{ vectorScalar }}(vec4<f32>(mish_value(fv.x), mish_value(fv.y), mish_value(fv.z), mish_value(fv.w)));
 }

+// Loads and stores vec4<T> while evaluating the selected unary operation per
+// component.
 {{ env.wgsl.resourceDeclarations }}
+{% macro emit_tanh_safe() %}
 fn tanh_safe(x: f32) -> f32 {
+  // tanh rounds to its saturated value for these tails in f32. Return that
+  // value directly, including for infinite input, before invoking the builtin.
   if (x > 10.0) { return 1.0; }
   if (x < -10.0) { return -1.0; }
+  // For tiny |x|, return x directly to preserve its sign and magnitude without
+  // relying on backend-specific builtin behavior near zero.
   if (x > -1.0e-4 && x < 1.0e-4) { return x; }
   return tanh(x);
+}{% endmacro %}
+{{ emit_tanh_safe() }}
 fn mish_log1p_small_positive(y: f32) -> f32 {
   // log1p(y) via series for small positive y, avoiding the log(1.0 + y)
+  // cancellation where 1.0 + y rounds back to 1.0.
   return y * (1.0 + y * (-0.5 + y * (0.3333333333333333 + y * (-0.25 + y * 0.2))));
 }
 fn mish_softplus(v: f32) -> f32 {
   // Numerically stable softplus = log(1 + e^v). The naive log(1.0 + exp(v))
   // collapses to 0 once exp(v) underflows below the 1.0 + x rounding step
+  // (v <~ -16), which zeros Mish's small finite negative tail. Use a series for
+  // the tiny tail and the identity-shifted logarithm otherwise.
   if (v > 20.0) { return v; }
   if (v < -8.0) { return mish_log1p_small_positive(exp(v)); }
   return log(1.0 + exp(-abs(v))) + max(v, 0.0);
   return v * tanh_safe(softplus);
 }
+{% set vec4PerThread = vec4PerThread %}
+{% if vec4PerThread > 1 %}
+const ITEMS: u32 = {{ vec4PerThread }}u;
+{% endif %}
 @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
+  // per-axis dispatch fold width (the dispatch caps x and spills the rest into y).
+{% if vec4PerThread > 1 %}
+  // Each invocation walks ITEMS vec4 groups a span apart. Consecutive lanes
+  // access consecutive words on every step, while each lane can keep several
+  // independent loads in flight.
+  let tid = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
+  let span = (params.count + ITEMS - 1u) / ITEMS;
+  for (var j = 0u; j < ITEMS; j = j + 1u) {
+    let i = tid + j * span;
+    if (i >= params.count) {
+      break;
+    }
+{% else %}
+  let i = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   if (i >= params.count) {
     return;
   }
+{% endif %}
   let xv = x[i];
   let fv = vec4<f32>(xv);
   y[i] = {{ vectorScalar }}(vec4<f32>(mish_value(fv.x), mish_value(fv.y), mish_value(fv.z), mish_value(fv.w)));
+{% if vec4PerThread > 1 %}
+  }
+{% endif %}
 }