sync 91d990483a17

Browse files

Files changed (8) hide show

README.md +11 -7
build/webgpu/bench.json +0 -1
build/webgpu/detect-neither.wgsl.jinja +1 -2
build/webgpu/manifest.json +55 -90
build/webgpu/metadata.json +16 -9
build/webgpu/test.json +9 -6
build/webgpu/unary-scalar.wgsl.jinja +8 -7
build/webgpu/unary-vec4.wgsl.jinja +29 -10

README.md CHANGED Viewed

@@ -18,15 +18,15 @@ See the [ONNX `IsInf` spec](https://onnx.ai/onnx/operators/onnx__IsInf.html) for
 ## Inputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `X` | `x` | `T` | — | — | Input floating-point tensor. | required |
 ## Outputs
-| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
-| `Y` | `y` | `B` | same as `X` | same as `X` | Boolean output tensor with the same shape as the input; `true` where the input is infinite. | required |
 ## Attributes
@@ -46,7 +46,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
@@ -56,10 +56,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` | — | — | Input floating-point tensor. | required |
 ## Outputs
+| Name | Upstream name | Logical dtype | Rank | Shape | Description | Presence |
 | --- | --- | --- | --- | --- | --- | --- |
+| `y` | `Y` | `B` | same as `x` | same as `x` | Boolean output tensor with the same shape as the input; `true` where the input is infinite. | 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.IsInf",
   "cases": [
     {
       "name": "isinf-f32-1m",

 {
   "cases": [
     {
       "name": "isinf-f32-1m",

build/webgpu/detect-neither.wgsl.jinja CHANGED Viewed

@@ -3,11 +3,10 @@
 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(
   @builtin(global_invocation_id) gid: vec3<u32>,
-  @builtin(num_workgroups) nwg: vec3<u32>,
 ) {
   // Four scalar stores retain arbitrary output lengths while sharing the
   // folded dispatch index and bounds work.
-  let item = gid.x + gid.y * nwg.x * {{ tunables.WORKGROUP_SIZE }}u;
   let base = item * 4u;
   if (base >= params.count) {
     return;

 @compute @workgroup_size({{ tunables.WORKGROUP_SIZE }})
 fn main(
   @builtin(global_invocation_id) gid: vec3<u32>,
 ) {
   // Four scalar stores retain arbitrary output lengths while sharing the
   // folded dispatch index and bounds work.
+  let item = gid.x + gid.y * {{ DISPATCH_FOLD_WIDTH }}u * {{ tunables.WORKGROUP_SIZE }}u;
   let base = item * 4u;
   if (base >= params.count) {
     return;

build/webgpu/manifest.json CHANGED Viewed

@@ -2,127 +2,92 @@
   "domain": "ai.onnx",
   "name": "IsInf",
   "sinceVersion": 20,
-  "description": "Maps each element of a floating-point tensor to `true` if it is infinite and `false` otherwise. The `detect_negative` and `detect_positive` attributes independently control whether negative and positive infinity are considered infinite for this mapping.",
-  "inputs": [{ "role": "X", "dtype": "T", "description": "Input floating-point tensor." }],
-  "outputs": [
-    {
-      "role": "Y",
-      "dtype": "B",
-      "rank": "ranks.X",
-      "description": "Boolean output tensor with the same shape as the input; `true` where the input is infinite.",
-      "shape": "shapes.X"
-    }
-  ],
-  "attributes": { "detect_negative": 1, "detect_positive": 1 },
-  "attributeDescriptions": {
-    "detect_negative": "When 1 (default), negative infinity maps to true; set to 0 to map negative infinity to false.",
-    "detect_positive": "When 1 (default), positive infinity maps to true; set to 0 to map positive infinity to false."
-  },
   "attributeConstraints": { "detect_negative": { "values": [0, 1] }, "detect_positive": { "values": [0, 1] } },
   "typeConstraints": { "T": ["float32", "float16"], "B": ["bool"] },
-  "args": {
-    "x": { "kind": "tensor", "semantic": "X", "role": "input" },
-    "y": { "kind": "tensor", "semantic": "Y", "role": "output" }
-  },
-  "tunables": { "WORKGROUP_SIZE": 256 },
-  "bindingSets": {
-    "sameLayoutVec4": [
-      {
-        "name": "x",
-        "arg": "x",
-        "semantic": "X",
-        "buffer": { "type": "read-only-storage" },
-        "elementType": "$vectorScalar"
-      },
-      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$outVec" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y) / 4" }] }
-      }
-    ],
-    "detectNeither": [
-      { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "u32" },
-      {
-        "name": "params",
-        "semantic": "kernel.params",
-        "buffer": { "type": "uniform" },
-        "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }] }
-      }
-    ]
-  },
-  "derive": {},
-  "constants": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
   "variants": [
     {
       "id": "detect_neither",
-      "when": ["attrs.detect_negative == 0", "attrs.detect_positive == 0", "numel(shapes.x) == numel(shapes.y)", "ranks.y == ranks.x"],
       "passes": [
         {
           "id": "main",
           "name": "IsInf.detectNeither",
-          "source": { "shader": "detect-neither.wgsl.jinja", "inputs": {} },
-          "bindings": "detectNeither",
-          "dispatch": { "threads": "ceilDiv(numel(shapes.y), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
-      ],
-      "priority": 30
     },
     {
       "id": "same_layout_vec4",
-      "when": ["attrs.detect_negative == 1 or attrs.detect_positive == 1", "numel(shapes.x) > 0", "numel(shapes.x) % 4 == 0", "numel(shapes.x) == numel(shapes.y)", "ranks.y == ranks.x", "f16Ok(dtypes.T)"],
-      "constants": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"", "outVec": "\"vec4<\" ~ dtypes.B ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "IsInf.vec4",
-          "source": {
-            "shader": "unary-vec4.wgsl.jinja",
-            "inputs": {
-              "op": "\"isinf\"",
-              "detectNegative": "attrs.detect_negative",
-              "detectPositive": "attrs.detect_positive"
-            }
           },
-          "bindings": "sameLayoutVec4",
-          "dispatch": { "threads": "numel(shapes.y) / 4", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
-      ],
-      "priority": 20
     },
     {
       "id": "elementwise",
-      "when": ["ranks.y == ranks.x", "numel(shapes.x) == numel(shapes.y)", "f16Ok(dtypes.T)"],
       "passes": [
         {
           "id": "main",
           "name": "IsInf",
-          "source": { "shader": "unary-scalar.wgsl.jinja", "inputs": { "op": "\"isinf\"", "itemsPerInvocation": 4 } },
           "bindings": [
-            {
-              "name": "x",
-              "arg": "x",
-              "semantic": "X",
-              "buffer": { "type": "read-only-storage" },
-              "elementType": "$scalar"
-            },
-            { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "u32" },
             {
               "name": "params",
-              "semantic": "kernel.params",
-              "buffer": { "type": "uniform" },
-              "struct": {
-                "name": "Params",
-                "fields": [
-                  { "name": "count", "type": "u32", "value": "numel(shapes.y)" },
-                  { "name": "detectNegative", "type": "u32", "value": "attrs.detect_negative" },
-                  { "name": "detectPositive", "type": "u32", "value": "attrs.detect_positive" }
-                ]
-              }
             }
           ],
-          "dispatch": { "threads": "ceilDiv(numel(shapes.y), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
         }
       ]
     }

   "domain": "ai.onnx",
   "name": "IsInf",
   "sinceVersion": 20,
+  "inputs": { "x": { "onnx": "X", "dtype": "T" } },
+  "outputs": { "y": { "onnx": "Y", "dtype": "B", "rank": "ranks.x", "shape": "shapes.x" } },
+  "attributes": { "detect_negative": { "default": 1 }, "detect_positive": { "default": 1 } },
   "attributeConstraints": { "detect_negative": { "values": [0, 1] }, "detect_positive": { "values": [0, 1] } },
   "typeConstraints": { "T": ["float32", "float16"], "B": ["bool"] },
+  "tunables": { "WORKGROUP_SIZE": { "default": 256 } },
+  "derive": { "scalar": "dtypes.T", "usesF16": "dtypes.T == \"f16\"" },
+  "when": ["numel(shapes.x) == numel(shapes.y)", "ranks.y == ranks.x"],
+  "bindings": { "y": { "buffer": "storage", "elementType": "u32" } },
   "variants": [
     {
       "id": "detect_neither",
+      "priority": 30,
+      "when": ["attrs.detect_negative == 0", "attrs.detect_positive == 0"],
       "passes": [
         {
           "id": "main",
           "name": "IsInf.detectNeither",
+          "shader": "detect-neither.wgsl.jinja",
+          "bindings": [
+            "y",
+            { "name": "params", "struct": [{ "name": "count", "type": "u32", "value": "numel(shapes.y)" }] }
+          ],
+          "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",
+      "priority": 20,
+      "when": ["attrs.detect_negative == 1 or attrs.detect_positive == 1", "numel(shapes.x) > 0", "numel(shapes.x) % 4 == 0", "f16Ok(dtypes.T)"],
+      "derive": { "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\"", "outVec": "\"vec4<\" ~ dtypes.B ~ \">\"" },
       "passes": [
         {
           "id": "main",
           "name": "IsInf.vec4",
+          "shader": "unary-vec4.wgsl.jinja",
+          "derive": {
+            "op": "\"isinf\"",
+            "detectNegative": "attrs.detect_negative",
+            "detectPositive": "attrs.detect_positive",
+            "vec4PerThread": "4 if numel(shapes.y) * dtypeBytes(tensorDtypes.y) <= 16777216 else 1"
           },
+          "bindings": [
+            { "arg": "x", "elementType": "$vectorScalar" },
+            { "arg": "y", "elementType": "$outVec" },
+            { "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": ["f16Ok(dtypes.T)"],
       "passes": [
         {
           "id": "main",
           "name": "IsInf",
+          "shader": "unary-scalar.wgsl.jinja",
+          "derive": { "op": "\"isinf\"", "itemsPerInvocation": 4 },
           "bindings": [
+            "x",
+            "y",
             {
               "name": "params",
+              "struct": [
+                { "name": "count", "type": "u32", "value": "numel(shapes.y)" },
+                { "name": "detectNegative", "type": "u32", "value": "attrs.detect_negative" },
+                { "name": "detectPositive", "type": "u32", "value": "attrs.detect_positive" }
+              ]
             }
           ],
+          "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,20 +1,27 @@
 {
   "name": "ai.onnx.IsInf",
-  "id": "_ai_onnx_isinf_webgpu_f721f3c",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "bench.json": "fpWJYObMRn5yexPegoVWy7o9/p0savu4NWVgbY8dPd4=",
-      "detect-neither.wgsl.jinja": "Og33gAC0ymZagfIJCnVu4N+uU8W0C+mq+mYg/9P12MQ=",
-      "manifest.json": "rPgpA/3IjGWEs8AY4VCBvJ4+d4Pe9BWpJ3X2RCCHERU=",
-      "test.json": "cnx4XDeVNUWmKXbZkFQuTOF/JuVvDXCK9zR+JzjXcI0=",
-      "unary-scalar.wgsl.jinja": "dewu+DcCE6Ec/dQBnj0eJC/5VIhhRWfvXJbqzZP3mGM=",
-      "unary-vec4.wgsl.jinja": "s46gUeMNEeKHdLt4//Eb8XKW6lIJU+ZHR9wj1nxqcIM="
     }
   },
-  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
-  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.IsInf" }
 }

 {
   "name": "ai.onnx.IsInf",
+  "id": "_ai_onnx_isinf_webgpu_a2a84e4",
   "version": 1,
   "license": "Apache-2.0",
   "backend": { "type": "webgpu" },
   "digest": {
     "algorithm": "sha256",
     "files": {
+      "bench.json": "GwHMLvIe56WI1ZJMsKpWAk+Q1myS7TBHG8aiVFxE3j4=",
+      "detect-neither.wgsl.jinja": "0JrGKclRlUw89VVRMFVfRSZ9hCrh2srIwSTAT48xjAM=",
+      "manifest.json": "PKpOiTfopy5MSaM8l+xoMZsdeZceGA/cpV917UyQgoE=",
+      "test.json": "MQmJcYVSPCKAMfiYgHGwSLwqUY7KumLrp7QsKeEIrqk=",
+      "unary-scalar.wgsl.jinja": "5Gu9hbmocQhPBlp4Bqc0XqwgbHsPeNN1GkKzDdSxcT4=",
+      "unary-vec4.wgsl.jinja": "1eB8LyA7gHU47DxSGRS/piI/bv4otnBUGTRnUi1X9tI="
     }
   },
+  "provenance": { "kernel": { "sha": "91d990483a174128daf7673f3f37a7c890493ae1", "dirty": false } },
+  "webgpu": {
+    "manifestSpec": "2.0",
+    "variants": {
+      "detect_neither": ["detect-neither.wgsl.jinja"],
+      "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.IsInf",
   "cases": [
     {
       "name": "f32_detect_both",
@@ -46,7 +45,7 @@
         "x": {
           "dtype": "float32",
           "shape": [5],
-          "data": { "kind": "values", "values": ["Infinity", "-Infinity", "NaN", 1.0, -1.0] }
         }
       },
       "outputs": { "y": { "dtype": "bool", "shape": [5], "tolerance": 0 } }
@@ -279,7 +278,7 @@
           "shape": [8],
           "data": {
             "kind": "values",
-            "values": ["Infinity", "-Infinity", "NaN", 0.0, -1.5, 3.6, "Infinity", "-Infinity"]
           }
         }
       },
@@ -337,10 +336,14 @@
     {
       "name": "f16_max_finite_not_inf",
       "inputs": {
-        "x": { "dtype": "float16", "shape": [4], "data": { "kind": "values", "values": [65504.0, -65504.0, 0.0, 1.0] } }
       },
       "outputs": {
-        "y": { "dtype": "bool", "shape": [4], "tolerance": 0, "data": { "kind": "values", "values": [0, 0, 0, 0] } }
       }
     },
     {
@@ -357,7 +360,7 @@
       },
       "outputs": { "y": { "dtype": "bool", "shape": [1, 1, 1, 1, 1, 1, 8], "tolerance": 0 } },
       "provenance": {
-        "notes": "Rank-7 flat input on the shared vec4 unary kernel (indexes by flat numel, so rank is irrelevant; sibling Cast/LeakyRelu impose no rank cap). Covers the regime the removed 'ranks.x <= 6' guard rejected."
       }
     },
     {

 {
   "cases": [
     {
       "name": "f32_detect_both",
         "x": {
           "dtype": "float32",
           "shape": [5],
+          "data": { "kind": "values", "values": ["Infinity", "Infinity", "Infinity", "Infinity", "Infinity"] }
         }
       },
       "outputs": { "y": { "dtype": "bool", "shape": [5], "tolerance": 0 } }
           "shape": [8],
           "data": {
             "kind": "values",
+            "values": ["Infinity", "Infinity", "Infinity", "Infinity", "Infinity", "Infinity", "Infinity", "Infinity"]
           }
         }
       },
     {
       "name": "f16_max_finite_not_inf",
       "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [4],
+          "data": { "kind": "values", "values": [65504.0, -65504.0, "Infinity", 1.0] }
+        }
       },
       "outputs": {
+        "y": { "dtype": "bool", "shape": [4], "tolerance": 0, "data": { "kind": "values", "values": [0, 0, 1, 0] } }
       }
     },
     {
       },
       "outputs": { "y": { "dtype": "bool", "shape": [1, 1, 1, 1, 1, 1, 8], "tolerance": 0 } },
       "provenance": {
+        "notes": "A rank-7 flat input exercises the shared vec4 unary path, whose indexing depends only on total element count."
       }
     },
     {

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

@@ -1,21 +1,22 @@
 {% 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 %}

 {% 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 %}

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

@@ -1,19 +1,35 @@
-// 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 }}
 @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);
   let bits = bitcast<vec4<u32>>(fv);
@@ -21,15 +37,18 @@ fn main(@builtin(global_invocation_id) gid: vec3<u32>, @builtin(num_workgroups)
   let mantissa = bits & vec4<u32>(0x007fffffu);
   let sgn = bits & vec4<u32>(0x80000000u);
   let is_inf = (exponent == vec4<u32>(0x7f800000u)) & (mantissa == vec4<u32>(0u));
-{% if source.detectPositive is defined and source.detectPositive %}
   let positive = is_inf & (sgn == vec4<u32>(0u));
 {% else %}
   let positive = vec4<bool>(false);
 {% endif %}
-{% if source.detectNegative is defined and source.detectNegative %}
   let negative = is_inf & (sgn != vec4<u32>(0u));
 {% else %}
   let negative = vec4<bool>(false);
 {% endif %}
   y[i] = select(vec4<u32>(0u), vec4<u32>(1u), positive | negative);
 }

+// Loads and stores vec4<T> while evaluating the selected unary operation per
+// component.
 {{ env.wgsl.resourceDeclarations }}
+{% 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);
   let bits = bitcast<vec4<u32>>(fv);
   let mantissa = bits & vec4<u32>(0x007fffffu);
   let sgn = bits & vec4<u32>(0x80000000u);
   let is_inf = (exponent == vec4<u32>(0x7f800000u)) & (mantissa == vec4<u32>(0u));
+{% if detectPositive is defined and detectPositive %}
   let positive = is_inf & (sgn == vec4<u32>(0u));
 {% else %}
   let positive = vec4<bool>(false);
 {% endif %}
+{% if detectNegative is defined and detectNegative %}
   let negative = is_inf & (sgn != vec4<u32>(0u));
 {% else %}
   let negative = vec4<bool>(false);
 {% endif %}
   y[i] = select(vec4<u32>(0u), vec4<u32>(1u), positive | negative);
+{% if vec4PerThread > 1 %}
+  }
+{% endif %}
 }