sync 2e7068faf55e

Browse files

Files changed (7) hide show

README.md +57 -0
build/webgpu/bench.json +41 -0
build/webgpu/manifest.json +80 -0
build/webgpu/metadata.json +19 -0
build/webgpu/test.json +210 -0
build/webgpu/unary-scalar.wgsl.jinja +44 -0
build/webgpu/unary-vec4.wgsl.jinja +46 -0

README.md CHANGED Viewed

@@ -1,3 +1,60 @@
 ---
 license: apache-2.0
 ---

 ---
+library_name: kernels
 license: apache-2.0
+tags:
+- kernel
+- webgpu
+- wgsl
 ---
+# ai.onnx.Relu
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 14
+## Description
+Applies the rectified linear unit function elementwise: `y = max(0, x)`. The output has the same shape and type as the input.
+See the [ONNX `Relu` spec](https://onnx.ai/onnx/operators/onnx__Relu.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `X` | `x` | `T` | — | — | Values clamped elementwise to a minimum of zero. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `Y` | `y` | `T` | same as `X` | same as `X` | Output tensor; same shape as the input. | required |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16`, `int32`, `int16`, `int8` |
+## 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
+- [`unary-scalar.wgsl.jinja`](build/webgpu/unary-scalar.wgsl.jinja)
+- [`unary-vec4.wgsl.jinja`](build/webgpu/unary-vec4.wgsl.jinja)
+## 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.
+```js
+import { getKernel } from "@huggingface/kernels";
+const kernel = await getKernel("webgpu-kernels/ai.onnx.Relu", { version: 1 });
+const { y } = await kernel({ x: { data: xData, shape: [] } });
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,41 @@

+{
+  "op": "ai.onnx.Relu",
+  "tunableSpace": { "WORKGROUP_SIZE": [64, 128, 256] },
+  "cases": [
+    {
+      "name": "relu-f32-1m",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "count": 1048576 },
+      "inputs": { "x": { "shape": [1048576], "dtype": "float32", "dist": "normal", "seed": 101, "scale": 2 } },
+      "outputs": { "y": { "shape": [1048576], "dtype": "float32" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }]
+      }
+    },
+    {
+      "name": "relu-f32-8m",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "count": 8388608 },
+      "inputs": { "x": { "shape": [8388608], "dtype": "float32", "seed": 7015, "dist": "normal", "scale": 2 } },
+      "outputs": { "y": { "shape": [8388608], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "relu-f16-8m",
+      "preset": "smoke",
+      "vars": { "dtype": "float16", "count": 8388608 },
+      "inputs": { "x": { "shape": [8388608], "dtype": "float16", "seed": 7016, "dist": "normal", "scale": 2 } },
+      "outputs": { "y": { "shape": [8388608], "dtype": "float16" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "relu-f32-scalar-fallback-8m",
+      "preset": "edge",
+      "vars": { "dtype": "float32", "count": 8388607 },
+      "inputs": { "x": { "shape": [8388607], "dtype": "float32", "dist": "normal", "seed": 7025, "scale": 2 } },
+      "outputs": { "y": { "shape": [8388607], "dtype": "float32", "dist": "empty" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,80 @@

+{
+  "domain": "ai.onnx",
+  "name": "Relu",
+  "sinceVersion": 14,
+  "description": "Applies the rectified linear unit function elementwise: `y = max(0, x)`. The output has the same shape and type as the input.",
+  "inputs": [{ "role": "X", "dtype": "T", "description": "Values clamped elementwise to a minimum of zero." }],
+  "outputs": [
+    {
+      "role": "Y",
+      "dtype": "T",
+      "rank": "ranks.x",
+      "shape": "shapes.x",
+      "description": "Output tensor; same shape as the input."
+    }
+  ],
+  "typeConstraints": { "T": ["float32", "float16", "int32", "int16", "int8"] },
+  "args": {
+    "x": { "kind": "tensor", "semantic": "X", "role": "input" },
+    "y": { "kind": "tensor", "semantic": "Y", "role": "output" }
+  },
+  "tunables": { "WORKGROUP_SIZE": 256 },
+  "derive": {
+    "wideVec4StorageOk": "dtypes.T != \"f32\" or device.features.has(\"subgroups\") or not (has(device.adapterInfo, \"subgroupMinSize\") and has(device.adapterInfo, \"subgroupMaxSize\") and device.adapterInfo.subgroupMinSize == 8 and device.adapterInfo.subgroupMaxSize <= 32)"
+  },
+  "variants": [
+    {
+      "id": "same_layout_vec4",
+      "when": ["numel(shapes.X) > 0", "numel(shapes.X) % 4 == 0", "numel(shapes.X) == numel(shapes.Y)", "f16Ok(dtypes.T)", "wideVec4StorageOk"],
+      "passes": [
+        {
+          "id": "main",
+          "name": "Relu.vec4",
+          "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" },
+          "source": { "shader": "unary-vec4.wgsl.jinja", "inputs": { "op": "\"relu\"" } }
+        }
+      ],
+      "priority": 20,
+      "constants": {
+        "scalar": "dtypes.T",
+        "usesF16": "dtypes.T == \"f16\"",
+        "vectorScalar": "\"vec4<\" ~ dtypes.T ~ \">\""
+      }
+    },
+    {
+      "id": "scalar",
+      "when": "f16Ok(dtypes.T)",
+      "passes": [
+        {
+          "id": "main",
+          "name": "scalar",
+          "source": { "shader": "unary-scalar.wgsl.jinja", "inputs": { "op": "\"relu\"", "itemsPerInvocation": 4 } },
+          "bindings": [
+            { "name": "x", "arg": "x", "semantic": "X", "buffer": { "type": "read-only-storage" }, "elementType": "$T" },
+            { "name": "y", "arg": "y", "semantic": "Y", "buffer": { "type": "storage" }, "elementType": "$T" },
+            {
+              "name": "params",
+              "semantic": "kernel.params",
+              "buffer": { "type": "uniform" },
+              "struct": { "name": "Params", "fields": [{ "name": "count", "type": "u32", "value": "numel(shapes.X)" }] }
+            }
+          ],
+          "dispatch": { "threads": "ceilDiv(numel(shapes.X), 4)", "workgroupSize": "tunables.WORKGROUP_SIZE" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

	@@ -0,0 +1,19 @@

+{
+  "name": "ai.onnx.Relu",
+  "id": "_ai_onnx_relu_webgpu_38d226d",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "pCIgxzDmyCADelXHsyZ3NqiYlep/O2kEemkYPRpHi4U=",
+      "manifest.json": "J17/rYpIwl28VO9DLnSBUimzXY15xRye4otbVNfgPOg=",
+      "test.json": "h0j4IDFTdi+LEdq6CjTJnSNQ1PTE9tpICFf8PVVFEKs=",
+      "unary-scalar.wgsl.jinja": "Vgm24ufHWZEC2cUeI0dA2snfsng2zIRpM1wC8URlqKA=",
+      "unary-vec4.wgsl.jinja": "xSW7Bah2pFe8MxojMV3phM8GEeicxBAk++vGxLURc0Q="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Relu" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,210 @@

+{
+  "op": "ai.onnx.Relu",
+  "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]
+  },
+  "cases": [
+    {
+      "name": "int16_scalar_boundaries",
+      "inputs": {
+        "x": { "dtype": "int16", "shape": [5], "data": { "kind": "values", "values": [-32768, -1, 0, 1, 32767] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "int16",
+          "shape": [5],
+          "tolerance": 0,
+          "data": { "kind": "values", "values": [0, 0, 0, 1, 32767] }
+        }
+      }
+    },
+    {
+      "name": "vector_17",
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [17], "data": { "kind": "fillFloat32", "sinStep": 0.37, "cosStep": 0.19 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [17], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "rank0_negative_scalar",
+      "inputs": { "x": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [-0.125] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "nan_input_propagates",
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": ["NaN", -1.0, 0.0, 2.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4], "tolerance": 0.000001, "allowNaN": true } }
+    },
+    {
+      "name": "f32_positive_subnormal_preserved_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: WebGPU/Metal flushes subnormals to zero (f32 and f16); the kernel cannot preserve denormal inputs/outputs bit-exactly."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/activation/activation_op_test.cc",
+        "test": "ActivationOpTest.Relu",
+        "notes": "Positive subnormal activations are valid Relu outputs; zero-flushing drops them while negative subnormals still clamp to zero."
+      },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [-1e-40, 0.0, 1e-40, 1e-39] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4], "tolerance": 0 } }
+    },
+    {
+      "name": "f32_positive_subnormal_preserved_scalar_gpu_gap",
+      "skipGpu": {
+        "category": "permanent",
+        "reason": "Portable WGSL floating-point semantics do not guarantee preservation of the subnormal values required by this fixture. Backend evidence: WebGPU/Metal flushes subnormals to zero (f32 and f16); the kernel cannot preserve denormal inputs/outputs bit-exactly."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/activation/activation_op_test.cc",
+        "test": "ActivationOpTest.Relu",
+        "notes": "Scalar-path companion: positive subnormal activations must pass through Relu unchanged."
+      },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-1e-40, 0.0, 1e-40] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3], "tolerance": 0 } }
+    },
+    {
+      "name": "int32_exact_above_float24",
+      "inputs": {
+        "x": {
+          "dtype": "int32",
+          "shape": [5],
+          "data": { "kind": "values", "values": [16777217, -16777217, 0, 123456789, -123456789] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [5] } }
+    },
+    {
+      "name": "ort_f32_activation_extremes",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/activation/activation_op_test.cc",
+        "test": "ActivationOpTest.Relu",
+        "notes": "Float32 subset of ORT's shared activation vector."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [13],
+          "data": {
+            "kind": "values",
+            "values": [-1.0, 0.0, 1.0, 100.0, -100.0, 1000.0, -1000.0, 1.1754943508222875e-38, 1.1754943508222876e-39, -1.1754943508222876e-39, 3.4028234663852886e+38, -3.4028234663852886e+38, "Infinity"]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [13], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_int8_activation_values",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/activation/activation_op_test.cc",
+        "test": "ActivationOpTest.Relu",
+        "notes": "Int8 values from ORT's Relu coverage."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [9],
+          "data": { "kind": "values", "values": [-1, -5, 0, 1, 5, 100, -100, -128, 127] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [9], "tolerance": 0 } }
+    },
+    {
+      "name": "ort_fp16_activation_extremes",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/activation/activation_op_test.cc",
+        "test": "ActivationOpTest.Relu_fp16"
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [13],
+          "data": {
+            "kind": "values",
+            "values": [-1.0, 0.0, 1.0, 100.0, -100.0, 1000.0, -1000.0, 1.1754943508222875e-38, 1.1754943508222876e-39, -1.1754943508222876e-39, 3.4028234663852886e+38, -3.4028234663852886e+38, "Infinity"]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [13], "tolerance": 0.001 } }
+    },
+    {
+      "name": "onnx_backend_rank3_float32",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_relu" },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4, 5], "tolerance": 0 } }
+    },
+    {
+      "name": "onnx_backend_relu",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_relu" },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": { "kind": "values", "values": { "$ref": "#/fixtureArrays/onnx_backend_input_x" } }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4, 5], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "vec4_f16_lanes",
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [16],
+          "data": {
+            "kind": "values",
+            "values": [-6.0, -4.0, -3.0, -2.0, -1.5, -1.0, -0.5, -0.25, 0.0, 0.25, 0.5, 1.0, 1.5, 2.0, 4.0, 6.0]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [16], "tolerance": 0 } }
+    },
+    {
+      "name": "vec4_i32_lanes",
+      "inputs": {
+        "x": { "dtype": "int32", "shape": [8], "data": { "kind": "values", "values": [-5, -1, 0, 3, 7, -100, 100, 2] } }
+      },
+      "outputs": { "y": { "dtype": "int32", "shape": [8] } }
+    },
+    {
+      "name": "vec4_i8_lanes",
+      "inputs": {
+        "x": {
+          "dtype": "int8",
+          "shape": [8],
+          "data": { "kind": "values", "values": [-128, -5, -1, 0, 1, 5, 100, 127] }
+        }
+      },
+      "outputs": { "y": { "dtype": "int8", "shape": [8] } }
+    },
+    {
+      "name": "empty_input_zero_dim",
+      "inputs": { "x": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [0], "tolerance": 0 } }
+    },
+    {
+      "name": "f16_nan_propagates_vec4",
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "values", "values": ["NaN", -2.0, -1.0, 0.0, 1.0, 2.0, -0.5, 0.5] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [8], "tolerance": 0, "allowNaN": true } },
+      "requires": { "features": ["shader-f16"] }
+    }
+  ]
+}

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

	@@ -0,0 +1,44 @@

+{% 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 T == "f16" %}
+enable f16;
+{% endif %}
+{{ env.wgsl.resourceDeclarations }}
+{% if T != "i32" %}
+fn is_nan_f32(value: f32) -> bool {
+  let bits = bitcast<u32>(value);
+  return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
+}
+{% endif %}
+{{ flat_tail_open() }}
+{% if T == "i32" %}
+  let value = x[i];
+  y[i] = select(value, 0i, value < 0i);
+{% else %}
+  let value = f32(x[i]);
+  var out = max(value, 0.0);
+  if (is_nan_f32(value)) {
+    out = value;
+  }
+  y[i] = {{ T }}(out);
+{% endif %}
+{{ flat_tail_close() -}}
+}

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

	@@ -0,0 +1,46 @@

+// 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 }}
+{% macro emit_is_nan_f32() %}
+fn is_nan_f32(value: f32) -> bool {
+  let bits = bitcast<u32>(value);
+  return (bits & 0x7f800000u) == 0x7f800000u && (bits & 0x007fffffu) != 0u;
+}{% endmacro %}
+{% if (source.op == "clip" and scalar != "i32" and scalar != "u32") or (source.op == "relu" and scalar != "i32") or source.op == "sin" or source.op == "tan" %}
+{{ emit_is_nan_f32() }}
+{% endif %}
+{% if source.op == "relu" and scalar != "i32" %}
+fn relu_value(value: f32) -> f32 {
+  // Some shader compilers apply no-NaN fast-math to vector select/compare
+  // expressions. Inspecting the IEEE payload explicitly keeps Relu's required
+  // NaN propagation deterministic across backends.
+  if (is_nan_f32(value)) {
+    return value;
+  }
+  return max(value, 0.0);
+}
+{% endif %}
+@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];
+{% if scalar == "i32" %}
+  y[i] = select(xv, vec4<i32>(0i), xv < vec4<i32>(0i));
+{% else %}
+  // Apply the IEEE NaN guard per lane because no-NaN fast math can rewrite
+  // vector compare/select.
+  let fv = vec4<f32>(xv);
+  y[i] = {{ vectorScalar }}(vec4<f32>(
+    relu_value(fv.x), relu_value(fv.y), relu_value(fv.z), relu_value(fv.w)));
+{% endif %}
+}