sync 91d990483a17

Browse files

Files changed (7) hide show

README.md +11 -7
build/webgpu/bench.json +0 -1
build/webgpu/manifest.json +37 -52
build/webgpu/metadata.json +11 -8
build/webgpu/test.json +3 -4
build/webgpu/unary-scalar.wgsl.jinja +11 -9
build/webgpu/unary-vec4.wgsl.jinja +31 -11

README.md CHANGED Viewed

@@ -18,15 +18,15 @@ See the [ONNX `Elu` spec](https://onnx.ai/onnx/operators/onnx__Elu.html) for the
 ## Inputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `X` | `x` | `T` | — | — | Values transformed elementwise by the ELU activation. | required |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `Y` | `y` | `T` | same as `X` | same as `X` | Output tensor with ELU applied elementwise; same shape as X. | required |
 ## Attributes
@@ -44,7 +44,7 @@ Default values (overridable per request):
 ## 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
@@ -53,10 +53,14 @@ Default values (overridable per request):
 ## 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 ELU activation. | required |
 ## Outputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `y` | `Y` | `T` | same as `x` | same as `x` | Output tensor with ELU applied elementwise; same shape as X. | required |
 ## Attributes
 ## 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.Elu",
   "cases": [
     {
       "name": "1m_f32",

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

build/webgpu/manifest.json CHANGED Viewed

@@ -2,81 +2,66 @@
   "domain": "ai.onnx",
   "name": "Elu",
   "sinceVersion": 6,
-  "description": "Applies the Exponential Linear Unit activation elementwise: `f(x) = x` for `x >= 0`, and `f(x) = alpha * (exp(x) - 1)` for `x < 0`. The `alpha` coefficient controls the saturation value for negative inputs.",
-  "inputs": [{ "role": "X", "dtype": "T", "description": "Values transformed elementwise by the ELU activation." }],
-  "outputs": [
-    {
-      "role": "Y",
-      "dtype": "T",
-      "rank": "ranks.X",
-      "description": "Output tensor with ELU applied elementwise; same shape as X.",
-      "shape": "shapes.X"
-    }
-  ],
-  "attributes": { "alpha": 1 },
-  "attributeDescriptions": {
-    "alpha": "Scaling coefficient applied to the exponential term for negative inputs; defaults to `1.0`."
-  },
   "typeConstraints": { "T": ["float32", "float16"] },
-  "args": {
-    "x": { "kind": "tensor", "semantic": "X", "role": "input" },
-    "y": { "kind": "tensor", "semantic": "Y", "role": "output" }
-  },
-  "tunables": { "WORKGROUP_SIZE": 256 },
-  "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
   "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": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "Elu.vec4",
-          "source": { "shader": "unary-vec4.wgsl.jinja", "inputs": { "op": "\"elu\"", "alpha": "attrs.alpha" } },
           "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) and (f16Ok(dtypes.T))",
       "passes": [
         {
           "id": "main",
           "name": "Elu",
-          "source": { "shader": "unary-scalar.wgsl.jinja", "inputs": { "op": "\"elu\"", "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)" },
-                  { "name": "alpha", "type": "f32", "value": "attrs.alpha" }
-                ]
-              }
             }
           ],
-          "dispatch": { "threads": "ceilDiv(numel(shapes.Y), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     }

   "domain": "ai.onnx",
   "name": "Elu",
   "sinceVersion": 6,
+  "inputs": { "x": { "onnx": "X", "dtype": "T" } },
+  "outputs": { "y": { "onnx": "Y", "dtype": "T", "rank": "ranks.x", "shape": "shapes.x" } },
+  "attributes": { "alpha": { "default": 1 } },
   "typeConstraints": { "T": ["float32", "float16"] },
+  "tunables": { "WORKGROUP_SIZE": { "default": 256 } },
+  "derive": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
   "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": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "Elu.vec4",
+          "shader": "unary-vec4.wgsl.jinja",
+          "derive": {
+            "op": "\"elu\"",
+            "alpha": "attrs.alpha",
+            "vec4PerThread": "4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1"
+          },
           "bindings": [
+            { "arg": "x", "elementType": "$vectorScalar" },
+            { "arg": "y", "elementType": "$vectorScalar" },
+            { "name": "params", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y) / 4" }] }
           ],
+          "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) and (f16Ok(dtypes.T))"],
       "passes": [
         {
           "id": "main",
           "name": "Elu",
+          "shader": "unary-scalar.wgsl.jinja",
+          "derive": { "op": "\"elu\"", "itemsPerInvocation": 4 },
           "bindings": [
+            "x",
+            "y",
             {
               "name": "params",
+              "struct": [
+                { "name": "count", "type": "u32", "value": "numel(shapes.y)" },
+                { "name": "alpha", "type": "f32", "value": "attrs.alpha" }
+              ]
             }
           ],
+          "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
+          }
         }
       ]
     }

build/webgpu/metadata.json CHANGED Viewed

@@ -1,19 +1,22 @@
 {
   "name": "ai.onnx.Elu",
-  "id": "_ai_onnx_elu_webgpu_e27eb6b",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "Fk685NfBNB1Zgf3q/ubBERo+pUczrEG/VeM9wSAtSok=",
-      "manifest.json": "JjDQbSvDifYG9piBDzIk8VmUoutDLm+HypEnzQnbcFE=",
-      "test.json": "PyjZAriKI7qWc4jh94V9Fm13hBYUuT9eK5Xlqeu9/24=",
-      "unary-scalar.wgsl.jinja": "jIfgc/qPA0E/qtbjHrzlkPXrKcr5NTj6MdZCGS4hk1g=",
-      "unary-vec4.wgsl.jinja": "NsfNFarG5uL16oqVeYJMOzArQ8At8c6otRbQZ697x6g="
     }
   },
-  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Elu" }
 }

 {
   "name": "ai.onnx.Elu",
+  "id": "_ai_onnx_elu_webgpu_ec52d7b",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "eiwKWjWWOzA7C9z26++rXtUz5yy870QjGR4DlTbyUxw=",
+      "manifest.json": "OzK0Gx6etOts1GvfOigK0vvpHTSzMjS+/QIH16Y2wmw=",
+      "test.json": "y6pMkqiL7rwwNdPkRUdOQj3ngy0de2tjjzkXYQirH7k=",
+      "unary-scalar.wgsl.jinja": "fBv4sZpdbUvk1gkgyMWwk82P4NXrnUMyuMwGhA4CXUg=",
+      "unary-vec4.wgsl.jinja": "JXyoLX0+h/n4xm23tuhcE47PvUrjjUk03Qo+Vuldn/s="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": { "same_layout_vec4": ["unary-vec4.wgsl.jinja"], "elementwise": ["unary-scalar.wgsl.jinja"] }
+  }
 }

build/webgpu/test.json CHANGED Viewed

@@ -1,5 +1,4 @@
 {
-  "op": "ai.onnx.Elu",
   "fixtureArrays": {
     "onnx_backend_input_x": [1.764052391052246, 0.40015721321105957, 0.978738009929657, 2.2408931255340576, 1.8675580024719238, -0.9772778749465942, 0.9500884413719177, -0.15135720372200012, -0.10321885347366333, 0.4105985164642334, 0.14404356479644775, 1.4542734622955322, 0.7610377073287964, 0.12167501449584961, 0.44386324286460876, 0.3336743414402008, 1.4940791130065918, -0.2051582634449005, 0.3130677044391632, -0.8540957570075989, -2.5529897212982178, 0.653618574142456, 0.8644362092018127, -0.7421650290489197, 2.269754648208618, -1.4543657302856445, 0.04575851559638977, -0.18718385696411133, 1.5327792167663574, 1.4693588018417358, 0.154947429895401, 0.37816253304481506, -0.8877857327461243, -1.980796456336975, -0.34791216254234314, 0.15634897351264954, 1.2302906513214111, 1.202379822731018, -0.38732680678367615, -0.302302747964859, -1.0485529899597168, -1.420017957687378, -1.7062702178955078, 1.950775384902954, -0.5096521973609924, -0.4380742907524109, -1.2527953386306763, 0.7774903774261475, -1.6138978004455566, -0.21274028718471527, -0.8954665660858154, 0.38690251111984253, -0.5108051300048828, -1.18063223361969, -0.02818222902715206, 0.4283318817615509, 0.06651721894741058, 0.30247190594673157, -0.6343221068382263, -0.3627411723136902]
   },
@@ -37,7 +36,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/activation/activation_op_test.cc",
         "test": "ActivationOpTest.Elu",
-        "notes": "ORT CPU rounds alpha * (exp(x) - 1) to zero for tiny negative inputs; current WebGPU leaks the original subnormal through."
       },
       "attrs": { "alpha": 0.5 },
       "inputs": {
@@ -75,7 +74,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/activation/activation_op_test.cc",
         "test": "ActivationOpTest.Elu",
-        "notes": "Scalar-path companion: tiny negative inputs round the expm1 branch to zero while positive subnormals pass through."
       },
       "attrs": { "alpha": 0.5 },
       "inputs": {
@@ -88,7 +87,7 @@
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/activation/activation_op_test.cc",
         "test": "ActivationOpTest.Elu",
-        "notes": "Scalar-path companion for the normal-range ELU cancellation edge."
       },
       "attrs": { "alpha": 0.1 },
       "inputs": {

 {
   "fixtureArrays": {
     "onnx_backend_input_x": [1.764052391052246, 0.40015721321105957, 0.978738009929657, 2.2408931255340576, 1.8675580024719238, -0.9772778749465942, 0.9500884413719177, -0.15135720372200012, -0.10321885347366333, 0.4105985164642334, 0.14404356479644775, 1.4542734622955322, 0.7610377073287964, 0.12167501449584961, 0.44386324286460876, 0.3336743414402008, 1.4940791130065918, -0.2051582634449005, 0.3130677044391632, -0.8540957570075989, -2.5529897212982178, 0.653618574142456, 0.8644362092018127, -0.7421650290489197, 2.269754648208618, -1.4543657302856445, 0.04575851559638977, -0.18718385696411133, 1.5327792167663574, 1.4693588018417358, 0.154947429895401, 0.37816253304481506, -0.8877857327461243, -1.980796456336975, -0.34791216254234314, 0.15634897351264954, 1.2302906513214111, 1.202379822731018, -0.38732680678367615, -0.302302747964859, -1.0485529899597168, -1.420017957687378, -1.7062702178955078, 1.950775384902954, -0.5096521973609924, -0.4380742907524109, -1.2527953386306763, 0.7774903774261475, -1.6138978004455566, -0.21274028718471527, -0.8954665660858154, 0.38690251111984253, -0.5108051300048828, -1.18063223361969, -0.02818222902715206, 0.4283318817615509, 0.06651721894741058, 0.30247190594673157, -0.6343221068382263, -0.3627411723136902]
   },
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/activation/activation_op_test.cc",
         "test": "ActivationOpTest.Elu",
+        "notes": "For tiny negative inputs, alpha*(exp(x)-1) rounds to zero; returning the original subnormal is incorrect."
       },
       "attrs": { "alpha": 0.5 },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/activation/activation_op_test.cc",
         "test": "ActivationOpTest.Elu",
+        "notes": "On an unaligned scalar path, tiny negative inputs round the expm1 branch to zero while positive subnormals pass through."
       },
       "attrs": { "alpha": 0.5 },
       "inputs": {
       "provenance": {
         "source": "onnxruntime/test/providers/cpu/activation/activation_op_test.cc",
         "test": "ActivationOpTest.Elu",
+        "notes": "Unaligned scalar-path inputs exercise ELU cancellation near zero in the normal range."
       },
       "attrs": { "alpha": 0.1 },
       "inputs": {

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

@@ -1,30 +1,32 @@
 {% 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 expm1_safe(x: f32) -> f32 {
   // exp(x) - 1 without the catastrophic cancellation of the literal exp(x) - 1.0,
   // which loses all significance for small |x| (exp(x) rounds to 1.0). A 6-term
-  // Maclaurin series is exact to < 1 ULP for |x| < 0.125 and matches the direct
-  // form at the switchover, so the normal range is unchanged.
   if (x > -0.125 && x < 0.125) {
     return x * (1.0 + x * (0.5 + x * (0.16666666666666666 + x * (0.041666666666666664 + x * (0.008333333333333333 + x * 0.001388888888888889)))));
   }

 {% 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 %}
+// Scalar unary elementwise implementation. Specialization emits only the
+// selected operation and any numerical helper it requires.
 {% if usesF16 %}
 enable f16;
 {% endif %}
 {{ env.wgsl.resourceDeclarations }}
+{% set alpha = alpha | default(0) %}
 fn expm1_safe(x: f32) -> f32 {
   // exp(x) - 1 without the catastrophic cancellation of the literal exp(x) - 1.0,
   // which loses all significance for small |x| (exp(x) rounds to 1.0). A 6-term
+  // Maclaurin series avoids that cancellation for |x| < 0.125; use the direct
+  // form elsewhere.
   if (x > -0.125 && x < 0.125) {
     return x * (1.0 + x * (0.5 + x * (0.16666666666666666 + x * (0.041666666666666664 + x * (0.008333333333333333 + x * 0.001388888888888889)))));
   }

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

@@ -1,30 +1,50 @@
-// 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 expm1_safe(x: vec4<f32>) -> vec4<f32> {
   // exp(x) - 1 without the catastrophic cancellation of the literal exp(x) - 1.0,
   // which loses all significance for small |x| (exp(x) rounds to 1.0). A 6-term
-  // Maclaurin series is exact to < 1 ULP for |x| < 0.125 and matches the direct
-  // form at the switchover, so the normal range is unchanged. Componentwise.
   let series = x * (1.0 + x * (0.5 + x * (0.16666666666666666 + x * (0.041666666666666664 + x * (0.008333333333333333 + x * 0.001388888888888889)))));
   let direct = exp(x) - 1.0;
   return select(direct, series, abs(x) < vec4<f32>(0.125));
 }
-const ELU_ALPHA: f32 = f32({{ source.alpha }});
 @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 }}(select(ELU_ALPHA * expm1_safe(fv), fv, fv >= vec4<f32>(0.0)));
 }

+// Loads and stores vec4<T> while evaluating the selected unary operation per
+// component.
 {{ env.wgsl.resourceDeclarations }}
+{% set alpha = alpha | default(0) %}
 fn expm1_safe(x: vec4<f32>) -> vec4<f32> {
   // exp(x) - 1 without the catastrophic cancellation of the literal exp(x) - 1.0,
   // which loses all significance for small |x| (exp(x) rounds to 1.0). A 6-term
+  // Maclaurin series avoids that cancellation for |x| < 0.125; use the direct
+  // form elsewhere. Componentwise.
   let series = x * (1.0 + x * (0.5 + x * (0.16666666666666666 + x * (0.041666666666666664 + x * (0.008333333333333333 + x * 0.001388888888888889)))));
   let direct = exp(x) - 1.0;
   return select(direct, series, abs(x) < vec4<f32>(0.125));
 }
+const ELU_ALPHA: f32 = f32({{ alpha }});
+{% 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 }}(select(ELU_ALPHA * expm1_safe(fv), fv, fv >= vec4<f32>(0.0)));
+{% if vec4PerThread > 1 %}
+  }
+{% endif %}
 }