sync 2e7068faf55e

Browse files

Files changed (7) hide show

README.md +57 -0
build/webgpu/bench.json +43 -0
build/webgpu/manifest.json +84 -0
build/webgpu/metadata.json +19 -0
build/webgpu/test.json +247 -0
build/webgpu/unary-scalar.wgsl.jinja +47 -0
build/webgpu/unary-vec4.wgsl.jinja +41 -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.Atanh
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 9
+## Description
+Computes the inverse hyperbolic tangent (`atanh`) of each element of the input tensor. Operates elementwise; input values must be in the open interval `(-1, 1)`.
+See the [ONNX `Atanh` spec](https://onnx.ai/onnx/operators/onnx__Atanh.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `input` | `x` | `T` | — | — | Values to map elementwise through inverse hyperbolic tangent; real-valued results require inputs in `(-1, 1)`. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `output` | `y` | `T` | same as `input` | same as `input` | The hyperbolic arctangent of each element of `input`; same shape as the input. | required |
+## Type constraints
+| Variable | Allowed dtypes |
+| --- | --- |
+| `T` | `float32`, `float16` |
+## 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.Atanh", { version: 1 });
+const { y } = await kernel({ x: { data: xData, shape: [] } });
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,43 @@

+{
+  "op": "ai.onnx.Atanh",
+  "cases": [
+    {
+      "name": "1m_f32",
+      "inputs": { "x": { "dtype": "float32", "shape": [1048576] } },
+      "outputs": { "y": { "dtype": "float32", "shape": [1048576] } }
+    },
+    {
+      "name": "1m_f16_vec4",
+      "preset": "smoke",
+      "vars": { "dtype": "float16", "count": 1048576 },
+      "inputs": {
+        "x": { "shape": [1048576], "dtype": "float16", "dist": "uniform", "seed": 911, "scale": 1.8, "offset": -0.9 }
+      },
+      "outputs": { "y": { "shape": [1048576], "dtype": "float16" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }]
+      }
+    },
+    {
+      "name": "1m_f32_scalar_fallback_unaligned",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "count": 1048575 },
+      "inputs": {
+        "x": { "shape": [1048575], "dtype": "float32", "dist": "uniform", "seed": 912, "scale": 1.8, "offset": -0.9 }
+      },
+      "outputs": { "y": { "shape": [1048575], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "1m_f16_scalar_fallback_unaligned",
+      "preset": "smoke",
+      "vars": { "dtype": "float16", "count": 1048575 },
+      "inputs": {
+        "x": { "shape": [1048575], "dtype": "float16", "dist": "uniform", "seed": 913, "scale": 1.8, "offset": -0.9 }
+      },
+      "outputs": { "y": { "shape": [1048575], "dtype": "float16" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    }
+  ]
+}

build/webgpu/manifest.json ADDED Viewed

	@@ -0,0 +1,84 @@

+{
+  "domain": "ai.onnx",
+  "name": "Atanh",
+  "sinceVersion": 9,
+  "description": "Computes the inverse hyperbolic tangent (`atanh`) of each element of the input tensor. Operates elementwise; input values must be in the open interval `(-1, 1)`.",
+  "inputs": [
+    {
+      "role": "input",
+      "dtype": "T",
+      "description": "Values to map elementwise through inverse hyperbolic tangent; real-valued results require inputs in `(-1, 1)`."
+    }
+  ],
+  "outputs": [
+    {
+      "role": "output",
+      "dtype": "T",
+      "rank": "ranks.input",
+      "description": "The hyperbolic arctangent of each element of `input`; same shape as the input.",
+      "shape": "shapes.input"
+    }
+  ],
+  "typeConstraints": { "T": ["float32", "float16"] },
+  "args": {
+    "x": { "kind": "tensor", "semantic": "input", "role": "input" },
+    "y": { "kind": "tensor", "semantic": "output", "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 ~ \">\""
+      },
+      "passes": [
+        {
+          "id": "main",
+          "name": "Atanh.vec4",
+          "source": { "shader": "unary-vec4.wgsl.jinja", "inputs": { "op": "\"atanh\"" } },
+          "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": "Atanh",
+          "source": { "shader": "unary-scalar.wgsl.jinja", "inputs": { "op": "\"atanh\"", "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" }
+        }
+      ]
+    }
+  ]
+}

build/webgpu/metadata.json ADDED Viewed

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

+{
+  "name": "ai.onnx.Atanh",
+  "id": "_ai_onnx_atanh_webgpu_8d63922",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "NTOrbyRKHuw4fvc5vl66WZf+8Jg8TH/srsgfC+k/shk=",
+      "manifest.json": "IC/2LTfUU97KTzRCcfiMS1d9aAETVID1Ch3F5dz2VA0=",
+      "test.json": "jBj/kfXXcx95j74iwm1vc6GhOfkFxhLPu2srA7Obz0U=",
+      "unary-scalar.wgsl.jinja": "k2knkcmGVcIrmqCemV/mIiIRauDUgvJBgfJmZjdeTbQ=",
+      "unary-vec4.wgsl.jinja": "2XlqwcUfyxUz90FaEDonfVXO1Pdijz2R2RJ+Mciv4Ks="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Atanh" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,247 @@

+{
+  "op": "ai.onnx.Atanh",
+  "cases": [
+    {
+      "name": "f32_values",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [-0.9, -0.5, 0.0, 0.25, 0.5, 0.9] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [6], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "f32_subnormal_identity_tail_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/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Atanh",
+        "notes": "For tiny finite inputs atanh(x) rounds back to x in float32; zero-flushing erases the signed tail."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [-1e-39, -1e-40, 1e-40, 1e-39] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4], "tolerance": 0 } }
+    },
+    {
+      "name": "f32_near_unit_boundary_accuracy_gpu_gap",
+      "skipGpu": {
+        "category": "todo",
+        "reason": "Near the high-slope domain boundary, the current f32 implementation differs from the CPU reference by about one ULP. A higher-accuracy f32 algorithm or software-extended precision could close this implementable gap."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Atanh",
+        "notes": "Values just inside the open Atanh domain are finite but numerically sensitive; this catches boundary accuracy loss before it becomes an Inf/NaN classification bug."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": {
+            "kind": "values",
+            "values": [-0.9999999403953552, -0.9999998807907104, 0.9999998807907104, 0.9999999403953552]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4], "tolerance": 5e-7, "relTolerance": 0 } }
+    },
+    {
+      "name": "f32_subnormal_identity_tail_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/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Atanh",
+        "notes": "Scalar-path companion: subnormal inputs are valid finite Atanh outputs."
+      },
+      "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": "f32_near_unit_boundary_accuracy_scalar_gpu_gap",
+      "skipGpu": {
+        "category": "todo",
+        "reason": "Near the high-slope domain boundary, the current scalar f32 implementation differs from the CPU reference by about one ULP. A higher-accuracy f32 algorithm or software-extended precision could close this implementable gap."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Atanh",
+        "notes": "Scalar-path companion for finite near-boundary Atanh accuracy just inside the open domain."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3],
+          "data": { "kind": "values", "values": [-0.9999999403953552, 0.9999998807907104, 0.9999999403953552] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3], "tolerance": 5e-7, "relTolerance": 0 } }
+    },
+    {
+      "name": "f16_values",
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 3],
+          "data": { "kind": "values", "values": [-0.9, -0.5, 0.0, 0.25, 0.5, 0.9] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 3] } },
+      "tolerance": 0.003
+    },
+    {
+      "name": "out_of_domain_returns_nan",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [1.5, -1.5, 2.0, -2.0, 1.0, -1.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [6], "tolerance": 0.000001, "allowNaN": true } }
+    },
+    {
+      "name": "f16_domain_and_nonfinite_values",
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [10],
+          "data": { "kind": "values", "values": [-1.0, -0.999, 0.0, 0.0, 0.999, 1.0, 1.001, -1.001, "Infinity", "NaN"] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [10], "tolerance": 0.003, "allowNaN": true } }
+    },
+    {
+      "name": "ort_float_boundary_interval",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Atanh"
+      },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [5], "data": { "kind": "values", "values": [-1.0, -0.5, 0.0, 0.5, 1.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [5], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "ort_domain_and_nonfinite_values",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Atanh",
+        "notes": "Adds endpoints, out-of-domain values, signed infinities, signed zero, and NaN."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [11],
+          "data": {
+            "kind": "values",
+            "values": [-1.0, -0.9999999, 0.0, 0.0, 0.9999999, 1.0, 1.0000001, -1.0000001, "Infinity", "-Infinity", "NaN"]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [11], "tolerance": 0.00001, "allowNaN": true } }
+    },
+    {
+      "name": "onnx_backend_example",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_atanh_example" },
+      "inputs": { "x": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-0.5, 0.0, 0.5] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "onnx_backend_atanh",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_atanh" },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [0.54881352186203, 0.7151893377304077, 0.6027633547782898, 0.5448831915855408, 0.42365479469299316, 0.6458941102027893, 0.4375872015953064, 0.891772985458374, 0.9636627435684204, 0.3834415078163147, 0.7917250394821167, 0.5288949012756348, 0.5680445432662964, 0.9255966544151306, 0.07103605568408966, 0.08712930232286453, 0.020218396559357643, 0.832619845867157, 0.7781567573547363, 0.8700121641159058, 0.978618323802948, 0.7991585731506348, 0.4614793658256531, 0.7805292010307312, 0.11827442795038223, 0.6399210095405579, 0.14335328340530396, 0.9446688890457153, 0.5218483209609985, 0.4146619439125061, 0.26455560326576233, 0.7742336988449097, 0.4561503231525421, 0.568433940410614, 0.018789799883961678, 0.6176354885101318, 0.6120957136154175, 0.6169340014457703, 0.9437480568885803, 0.681820273399353, 0.35950788855552673, 0.43703195452690125, 0.6976311802864075, 0.0602254718542099, 0.6667667031288147, 0.670637845993042, 0.21038256585597992, 0.12892629206180573, 0.31542834639549255, 0.36371076107025146, 0.5701967477798462, 0.4386015236377716, 0.9883738160133362, 0.10204481333494186, 0.20887675881385803, 0.16130951046943665, 0.6531082987785339, 0.25329160690307617, 0.4663107693195343, 0.24442559480667114]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4, 5], "tolerance": 0.00001, "allowNaN": true } }
+    },
+    {
+      "name": "onnx_backend_atanh_example",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_atanh_example" },
+      "inputs": { "x": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-0.5, 0.0, 0.5] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [3], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "vec4_f16_lanes",
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [16],
+          "data": {
+            "kind": "values",
+            "values": [-0.95, -0.9, -0.75, -0.5, -0.25, -0.125, 0.0, 0.0625, 0.125, 0.25, 0.375, 0.5, 0.625, 0.75, 0.875, 0.95]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [16], "tolerance": 0.001, "relTolerance": 0.005 } }
+    },
+    {
+      "name": "empty_input_zero_dim",
+      "inputs": { "x": { "dtype": "float32", "shape": [0], "data": { "kind": "values", "values": [] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [0], "tolerance": 0 } }
+    },
+    {
+      "name": "f32_near_unit_boundary_gpu_runnable",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [-0.9999, -0.999, -0.99, 0.99, 0.999, 0.9999] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [6], "tolerance": 0.001, "relTolerance": 0.0001 } }
+    },
+    {
+      "name": "rank0_scalar",
+      "inputs": { "x": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [0.5] } } },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [],
+          "data": { "kind": "values", "values": [0.5493061443340549] },
+          "tolerance": 0.000001
+        }
+      }
+    },
+    {
+      "name": "f32_boundary_exact_one_and_neg_one",
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [-1.0, -0.5, 0.5, 1.0] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4], "tolerance": 0.000001, "allowNaN": true } }
+    },
+    {
+      "name": "f16_empty_zero_numel_shape",
+      "inputs": { "x": { "dtype": "float16", "shape": [0], "data": { "kind": "values", "values": [] } } },
+      "outputs": { "y": { "dtype": "float16", "shape": [0] } }
+    },
+    {
+      "name": "f32_vec4_sustained_1024",
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [1024], "data": { "kind": "linspace", "start": -0.9, "end": 0.9 } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1024], "tolerance": 0.00001 } }
+    }
+  ]
+}

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

	@@ -0,0 +1,47 @@

+{% 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 }}
+{% set unaryDomainGuard = device.adapterInfo.architecture == "" or device.adapterInfo.architecture == "apple" %}
+{% macro emit_atanh_safe(guarded) %}
+fn atanh_safe(x: f32) -> f32 {
+{%- if guarded %}
+  // A no-NaN compilation path can return a finite value for out-of-domain
+  // operands. Classify the loaded bits instead: |x| == 1 is a signed
+  // infinity and |x| > 1 (including NaN) is a quiet NaN carrying the operand's
+  // sign.
+  let bits = bitcast<u32>(x);
+  let magnitude = bits & 0x7fffffffu;
+  if (magnitude == 0x3f800000u) {
+    return bitcast<f32>((bits & 0x80000000u) | 0x7f800000u);
+  }
+  if (magnitude > 0x3f800000u) {
+    return bitcast<f32>(bits | 0x7fc00000u);
+  }
+{%- endif %}
+  return atanh(x);
+}{% endmacro %}
+{{ emit_atanh_safe(unaryDomainGuard) }}
+{{ flat_tail_open() }}
+  let v = f32(x[i]);
+  y[i] = {{ scalar }}(atanh_safe(v));
+{{ flat_tail_close() -}}
+}

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

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

+// 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 }}
+{% set unaryDomainGuard = device.adapterInfo.architecture == "" or device.adapterInfo.architecture == "apple" %}
+{% macro emit_atanh_safe(guarded) %}
+fn atanh_safe(x: f32) -> f32 {
+{%- if guarded %}
+  // A no-NaN compilation path can return a finite value for out-of-domain
+  // operands. Classify the loaded bits instead: |x| == 1 is a signed
+  // infinity and |x| > 1 (including NaN) is a quiet NaN carrying the operand's
+  // sign.
+  let bits = bitcast<u32>(x);
+  let magnitude = bits & 0x7fffffffu;
+  if (magnitude == 0x3f800000u) {
+    return bitcast<f32>((bits & 0x80000000u) | 0x7f800000u);
+  }
+  if (magnitude > 0x3f800000u) {
+    return bitcast<f32>(bits | 0x7fc00000u);
+  }
+{%- endif %}
+  return atanh(x);
+}{% endmacro %}
+{{ emit_atanh_safe(unaryDomainGuard) }}
+@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>(
+    atanh_safe(fv.x), atanh_safe(fv.y), atanh_safe(fv.z), atanh_safe(fv.w)));
+}