sync 2e7068faf55e

Browse files

Files changed (7) hide show

README.md +57 -0
build/webgpu/bench.json +45 -0
build/webgpu/manifest.json +80 -0
build/webgpu/metadata.json +19 -0
build/webgpu/test.json +238 -0
build/webgpu/unary-scalar.wgsl.jinja +205 -0
build/webgpu/unary-vec4.wgsl.jinja +198 -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.Cos
+`ai.onnx`  ·  standard ONNX operator  ·  ONNX opset ≥ 7
+## Description
+Computes the cosine of each element in the input tensor. The output has the same shape and type as the input.
+See the [ONNX `Cos` spec](https://onnx.ai/onnx/operators/onnx__Cos.html) for the reference semantics.
+## Inputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `input` | `x` | `T` | — | — | Angles in radians whose cosine is computed elementwise. | required |
+## Outputs
+| Name | Bind key | Logical dtype | Rank | Shape | Description | Presence |
+| --- | --- | --- | --- | --- | --- | --- |
+| `output` | `y` | `T` | same as `input` | same as `input` | Elementwise cosine of the input tensor. | 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.Cos", { version: 1 });
+const { y } = await kernel({ x: { data: xData, shape: [] } });
+```

build/webgpu/bench.json ADDED Viewed

	@@ -0,0 +1,45 @@

+{
+  "op": "ai.onnx.Cos",
+  "cases": [
+    {
+      "name": "1m_f32",
+      "inputs": { "x": { "dtype": "float32", "shape": [1048576] } },
+      "outputs": { "y": { "dtype": "float32", "shape": [1048576] } }
+    },
+    {
+      "name": "cos-f32-1m-vec4-healthy",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "count": 1048576 },
+      "inputs": { "x": { "shape": [1048576], "dtype": "float32", "dist": "normal", "seed": 8801, "scale": 3 } },
+      "outputs": { "y": { "shape": [1048576], "dtype": "float32" } },
+      "bench": {
+        "primary": true,
+        "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }]
+      }
+    },
+    {
+      "name": "cos-f32-scalar-path-nonaligned",
+      "preset": "smoke",
+      "vars": { "dtype": "float32", "count": 1048577 },
+      "inputs": { "x": { "shape": [1048577], "dtype": "float32", "dist": "normal", "seed": 8802, "scale": 3 } },
+      "outputs": { "y": { "shape": [1048577], "dtype": "float32" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "cos-f16-1m-vec4",
+      "preset": "smoke",
+      "vars": { "dtype": "float16", "count": 1048576 },
+      "inputs": { "x": { "shape": [1048576], "dtype": "float16", "dist": "normal", "seed": 8803, "scale": 3 } },
+      "outputs": { "y": { "shape": [1048576], "dtype": "float16" } },
+      "bench": { "metrics": [{ "type": "bandwidth", "value": "args.count * dtypeBytes(args.dtype) * 2" }] }
+    },
+    {
+      "name": "cos-f16-scalar-path-nonaligned",
+      "preset": "smoke",
+      "vars": { "dtype": "float16", "count": 1048577 },
+      "inputs": { "x": { "shape": [1048577], "dtype": "float16", "dist": "normal", "seed": 8804, "scale": 3 } },
+      "outputs": { "y": { "shape": [1048577], "dtype": "float16" } },
+      "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": "Cos",
+  "sinceVersion": 7,
+  "description": "Computes the cosine of each element in the input tensor. The output has the same shape and type as the input.",
+  "inputs": [
+    { "role": "input", "dtype": "T", "description": "Angles in radians whose cosine is computed elementwise." }
+  ],
+  "outputs": [
+    {
+      "role": "output",
+      "dtype": "T",
+      "rank": "ranks.input",
+      "description": "Elementwise cosine of the input tensor.",
+      "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": "Cos.vec4",
+          "source": { "shader": "unary-vec4.wgsl.jinja", "inputs": { "op": "\"cos\"" } },
+          "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": "Cos",
+          "source": { "shader": "unary-scalar.wgsl.jinja", "inputs": { "op": "\"cos\"", "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.Cos",
+  "id": "_ai_onnx_cos_webgpu_cac3d94",
+  "version": 1,
+  "license": "Apache-2.0",
+  "backend": { "type": "webgpu" },
+  "digest": {
+    "algorithm": "sha256",
+    "files": {
+      "bench.json": "aMFVQ+6rhvrMngby96efMmsldwq0VRJ9P0ICGr91lUY=",
+      "manifest.json": "jHDCQwVeQCIpu3FV1PiYQD9ymkzPseA2K6yGVJf3wBk=",
+      "test.json": "4f0wcy6ce8pie62+H13dy87SkZOfWv4TkyeqalgROmw=",
+      "unary-scalar.wgsl.jinja": "VPyTY5cV6Zjidm4VpLmKEhg9VJDzG6N8+WZ/QJc+mlo=",
+      "unary-vec4.wgsl.jinja": "1g3ktdx4TPnyT48hkpcIvkPRV8APmLzL3NESPtJPryA="
+    }
+  },
+  "provenance": { "kernel": { "sha": "2e7068faf55e7f43df740015f6d1ee49391a41c5", "dirty": false } },
+  "webgpu": { "manifestSpec": "1.0", "specialized": true, "opPath": "ops/ai.onnx.Cos" }
+}

build/webgpu/test.json ADDED Viewed

	@@ -0,0 +1,238 @@

+{
+  "op": "ai.onnx.Cos",
+  "cases": [
+    {
+      "name": "f32_values",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [6],
+          "data": { "kind": "values", "values": [-3.0, -1.0, 0.0, 0.5, 1.0, 3.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [6], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "large_argument_range_reduction",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": {
+            "kind": "values",
+            "values": [1000000.0, 10000000.0, 10000000000000.0, 100000000000000000000.0, -1000000000000000.0, -123456.78, 314159.265, 2500000000.0]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [8], "tolerance": 0.0001 } }
+    },
+    {
+      "name": "f32_large_argument_range_reduction_accuracy_gpu_gap",
+      "skipGpu": {
+        "category": "todo",
+        "reason": "The current f32 large-argument range reduction differs from the CPU reference by about one ULP. A more accurate range reduction or software-extended precision could close this implementable gap."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.CosFloat",
+        "notes": "Finite large arguments stress Cos range reduction; this uses a stricter tolerance than the broad smoke case."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": {
+            "kind": "values",
+            "values": [1000000.0, 10000000.0, 10000000000000.0, 100000000000000000000.0, -1000000000000000.0, -123456.78, 314159.265, 2500000000.0]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [8], "tolerance": 1e-8, "relTolerance": 0 } }
+    },
+    {
+      "name": "f32_large_argument_range_reduction_accuracy_scalar_gpu_gap",
+      "skipGpu": {
+        "category": "todo",
+        "reason": "The current scalar f32 large-argument range reduction differs from the CPU reference by about one ULP. A more accurate range reduction or software-extended precision could close this implementable gap."
+      },
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.CosFloat",
+        "notes": "Scalar-path companion for strict large-argument Cos range-reduction accuracy."
+      },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [7],
+          "data": {
+            "kind": "values",
+            "values": [1000000.0, 10000000.0, 10000000000000.0, -1000000000000000.0, -123456.78, 314159.265, 2500000000.0]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [7], "tolerance": 1e-8, "relTolerance": 0 } }
+    },
+    {
+      "name": "special_values_infinity_nan",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [5],
+          "data": { "kind": "values", "values": ["-Infinity", 0.0, 0.0, "Infinity", "NaN"] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [5], "tolerance": 0.000001, "allowNaN": true } }
+    },
+    {
+      "name": "rank0_scalar",
+      "inputs": { "x": { "dtype": "float32", "shape": [], "data": { "kind": "values", "values": [1.25] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "ort_float_opset22",
+      "provenance": {
+        "source": "onnxruntime/test/providers/cpu/math/element_wise_ops_test.cc",
+        "test": "MathOpTest.Cos_Opset22"
+      },
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [4], "data": { "kind": "values", "values": [1.1, -1.1, 2.2, -2.2] } }
+      },
+      "outputs": {
+        "y": {
+          "dtype": "float32",
+          "shape": [4],
+          "tolerance": 0.000001,
+          "data": {
+            "kind": "values",
+            "values": [0.4535961151123047, 0.4535961151123047, -0.5885010957717896, -0.5885010957717896]
+          }
+        }
+      }
+    },
+    {
+      "name": "onnx_backend_example",
+      "provenance": {
+        "source": "cmake/external/onnx/onnx/backend/test/data/node/test_cos_example",
+        "test": "test_cos_example"
+      },
+      "inputs": { "x": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-1.0, 0.0, 1.0] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [3], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "onnx_backend_cos",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_cos" },
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [3, 4, 5],
+          "data": {
+            "kind": "values",
+            "values": [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]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [3, 4, 5], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "onnx_backend_cos_example",
+      "provenance": { "source": "cmake/external/onnx/onnx/backend/test/data/node/test_cos_example" },
+      "inputs": { "x": { "dtype": "float32", "shape": [3], "data": { "kind": "values", "values": [-1.0, 0.0, 1.0] } } },
+      "outputs": { "y": { "dtype": "float32", "shape": [3], "tolerance": 0.00001 } }
+    },
+    {
+      "name": "vec4_f32_nonfinite",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [8],
+          "data": { "kind": "values", "values": ["-Infinity", -1.0, 0.0, 0.5, 1.0, "Infinity", "NaN", 2.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [8], "tolerance": 0.000001, "allowNaN": true } }
+    },
+    {
+      "name": "vec4_f16_lanes",
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [16],
+          "data": {
+            "kind": "values",
+            "values": [-100.0, -50.0, -20.0, -10.0, -6.0, -3.0, -1.5, -0.5, 0.0, 0.5, 1.5, 3.0, 6.0, 10.0, 50.0, 100.0]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [16], "tolerance": 0.001, "relTolerance": 0.002 } }
+    },
+    {
+      "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_scalar_path_odd_numel",
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [7],
+          "data": { "kind": "values", "values": [-2.0, -1.0, -0.5, 0.0, 0.5, 1.0, 2.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [7], "tolerance": 0.001, "relTolerance": 0.002 } }
+    },
+    {
+      "name": "f16_realistic_2d_finite",
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [2, 64],
+          "data": {
+            "kind": "values",
+            "values": [-3.140625, -3.046875, -2.953125, -2.859375, -2.765625, -2.671875, -2.578125, -2.484375, -2.390625, -2.296875, -2.203125, -2.109375, -2.015625, -1.921875, -1.828125, -1.734375, -1.640625, -1.546875, -1.453125, -1.359375, -1.265625, -1.171875, -1.078125, -0.984375, -0.890625, -0.796875, -0.703125, -0.609375, -0.515625, -0.421875, -0.328125, -0.234375, -0.140625, -0.046875, 0.046875, 0.140625, 0.234375, 0.328125, 0.421875, 0.515625, 0.609375, 0.703125, 0.796875, 0.890625, 0.984375, 1.078125, 1.171875, 1.265625, 1.359375, 1.453125, 1.546875, 1.640625, 1.734375, 1.828125, 1.921875, 2.015625, 2.109375, 2.203125, 2.296875, 2.390625, 2.484375, 2.578125, 2.671875, 2.765625, -3.0, -2.5, -2.0, -1.5, -1.0, -0.5, 0.0, 0.5, 1.0, 1.5, 2.0, 2.5, 3.0, 0.25, 0.75, 1.25, 1.75, 2.25, 2.75, -0.25, -0.75, -1.25, -1.75, -2.25, -2.75, 0.125, 0.375, 0.625, 0.875, 1.125, 1.375, 1.625, 1.875, 2.125, 2.375, 2.625, 2.875, -0.125, -0.375, -0.625, -0.875, -1.125, -1.375, -1.625, -1.875, -2.125, -2.375, -2.625, -2.875, 3.0, -3.0, 1.0, -1.0, 0.5, -0.5, 2.0, -2.0, 1.5, -1.5, 2.5, -2.5, 0.0, 0.0, 3.0]
+          }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [2, 64], "tolerance": 0.001, "relTolerance": 0.002 } }
+    },
+    {
+      "name": "f32_scalar_path_finite_odd_numel",
+      "inputs": {
+        "x": { "dtype": "float32", "shape": [5], "data": { "kind": "values", "values": [-2.5, -0.75, 0.0, 1.25, 3.5] } }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [5], "tolerance": 0.000001 } }
+    },
+    {
+      "name": "reduce_threshold_boundary_10000",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [4],
+          "data": { "kind": "values", "values": [9999.0, 10000.0, -9999.0, -10000.0] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [4], "tolerance": 0.001 } }
+    },
+    {
+      "name": "f16_cos_range_finite_and_nonfinite_vec4",
+      "inputs": {
+        "x": {
+          "dtype": "float16",
+          "shape": [8],
+          "data": { "kind": "values", "values": ["-Infinity", -100.0, -1.0, 0.0, 1.0, 100.0, "Infinity", "NaN"] }
+        }
+      },
+      "outputs": { "y": { "dtype": "float16", "shape": [8], "tolerance": 0.002, "allowNaN": true } }
+    },
+    {
+      "name": "f32_vec4_sustained_1024",
+      "inputs": {
+        "x": {
+          "dtype": "float32",
+          "shape": [1024],
+          "data": { "kind": "linspace", "start": -6.283185, "end": 6.283185 }
+        }
+      },
+      "outputs": { "y": { "dtype": "float32", "shape": [1024], "tolerance": 0.00001 } }
+    }
+  ]
+}

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

	@@ -0,0 +1,205 @@

+{% 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 }}
+{% macro emit_reduce_pio2() %}
+fn reduce_pio2(ax: f32) -> Pio2 {
+  let ix = bitcast<u32>(ax);
+  let e = (ix >> 23u) & 0xffu;
+  let mant = (ix & 0x7fffffu) | 0x800000u;   // 24-bit mantissa
+  let p = i32(e) - 150;                        // ax = mant * 2^p
+  // 96-bit accumulator a2:a1:a0 holding (ax * 2/pi  mod 4) * 2^94.
+  var a0 = 0u; var a1 = 0u; var a2 = 0u;
+  for (var k = 0u; k < 14u; k = k + 1u) {
+    let shift = p - 24 * (i32(k) + 1) + 94;
+    if (shift <= -48 || shift >= 96) { continue; }
+    // prod = mant * TWO_OVER_PI[k]  (<= 48 bits) -> phi32:plo
+    let t = TWO_OVER_PI[k];
+    let mlo = mant & 0xffffu; let mhi = mant >> 16u;
+    let tlo = t & 0xffffu;    let thi = t >> 16u;
+    let ll = mlo * tlo;
+    let mid = mlo * thi + mhi * tlo;            // <= 25 bits, no overflow
+    let hh = mhi * thi;
+    let plo = ll + ((mid & 0xffffu) << 16u);
+    let carry = select(0u, 1u, plo < ll);
+    let phi32 = hh + (mid >> 16u) + carry;      // <= 16 bits
+    // Place (phi32:plo) << shift as a 64-bit value (v1:v0) at offset off >= 0.
+    var v0 = 0u; var v1 = 0u; var off = 0;
+    if (shift < 0) {
+      let s = u32(-shift);                      // 1..47
+      if (s < 32u) {
+        v0 = (plo >> s) | (phi32 << (32u - s));
+        v1 = phi32 >> s;
+      } else {
+        v0 = phi32 >> (s - 32u);
+        v1 = 0u;
+      }
+      off = 0;
+    } else {
+      v0 = plo; v1 = phi32; off = shift;        // 0..95
+    }
+    // Spread (v1:v0) across 96-bit limbs starting at bit `off`.
+    let li = off / 32; let b = u32(off % 32);
+    var w0 = 0u; var w1 = 0u; var w2 = 0u;
+    if (b == 0u) { w0 = v0; w1 = v1; w2 = 0u; }
+    else {
+      w0 = v0 << b;
+      w1 = (v1 << b) | (v0 >> (32u - b));
+      w2 = v1 >> (32u - b);
+    }
+    var t0 = 0u; var t1 = 0u; var t2 = 0u;
+    if (li == 0) { t0 = w0; t1 = w1; t2 = w2; }
+    else if (li == 1) { t1 = w0; t2 = w1; }
+    else { t2 = w0; }
+    // 96-bit add acc += t (bits above 2^96 are multiples of 4, so discarded).
+    let n0 = a0 + t0; let c0 = select(0u, 1u, n0 < a0);
+    let n1a = a1 + t1; let c1a = select(0u, 1u, n1a < a1);
+    let n1 = n1a + c0; let c1b = select(0u, 1u, n1 < n1a);
+    let n2 = a2 + t2 + c1a + c1b;
+    a0 = n0; a1 = n1; a2 = n2;
+  }
+  var n = a2 >> 30u;                            // integer part mod 4
+  var phi = f32(a2 & 0x3fffffffu) * (1.0 / 1073741824.0)
+          + f32(a1) * (1.0 / 4611686018427387904.0);   // fraction in [0,1)
+  if (phi >= 0.5) { phi = phi - 1.0; n = n + 1u; }
+  var out: Pio2;
+  out.octant = n & 3u;
+  out.r = phi * PIO2_F;
+  return out;
+}{% endmacro %}
+{% macro emit_reduce_pio2_fast() %}
+fn reduce_pio2_fast(ax: f32) -> Pio2 {
+  // Cody-Waite reduction for ordinary magnitudes. Splitting pi/2 keeps the
+  // residual accurate even when n is thousands; larger inputs use Payne-Hanek.
+  let nFloat = floor(fma(ax, INV_PIO2_F, 0.5));
+  var residual = fma(-nFloat, PIO2_HI_F, ax);
+  residual = fma(-nFloat, PIO2_LO_F, residual);
+  var out: Pio2;
+  out.octant = u32(nFloat) & 3u;
+  out.r = residual;
+  return out;
+}{% endmacro %}
+{% macro emit_trig_reduction_support() %}
+{% set needPreciseTrigCentered = false %}
+{% set needPreciseTrigTwoPi = false %}
+// Backend-stable f32 sine/cosine core.
+//
+// Shader transcendental accuracy is implementation-defined, and some portable
+// backends are only accurate to roughly 1e-4. Each path retains the most accurate
+// available phase representation, reduces it to [-pi, pi], then uses these
+// polynomials. The half-pi core is degree 13 for sine and degree 12 for cosine;
+// truncation error is well below one f32 ULP over its documented interval.
+{% if needPreciseTrigCentered %}
+const PRECISE_TRIG_PI: f32 = 3.141592653589793;
+{% endif %}
+{% if needPreciseTrigTwoPi %}
+const PRECISE_TRIG_TWO_PI: f32 = 6.283185307179586;
+{% endif %}
+{% if needPreciseTrigCentered %}
+const PRECISE_TRIG_HALF_PI: f32 = 1.5707963267948966;
+{% endif %}
+// Input must be in [-pi/2, pi/2]. Returns (cos(x), sin(x)).
+fn precise_sincos_half_pi(x: f32) -> vec2<f32> {
+  let x2 = x * x;
+  var sinPolynomial = 1.6059043836821613e-10;
+  sinPolynomial = fma(sinPolynomial, x2, -2.505210838544172e-8);
+  sinPolynomial = fma(sinPolynomial, x2, 2.7557319223985893e-6);
+  sinPolynomial = fma(sinPolynomial, x2, -1.984126984126984e-4);
+  sinPolynomial = fma(sinPolynomial, x2, 8.333333333333333e-3);
+  sinPolynomial = fma(sinPolynomial, x2, -1.6666666666666666e-1);
+  let sine = x * fma(sinPolynomial, x2, 1.0);
+  var cosPolynomial = 2.08767569878681e-9;
+  cosPolynomial = fma(cosPolynomial, x2, -2.755731922398589e-7);
+  cosPolynomial = fma(cosPolynomial, x2, 2.48015873015873e-5);
+  cosPolynomial = fma(cosPolynomial, x2, -1.388888888888889e-3);
+  cosPolynomial = fma(cosPolynomial, x2, 4.1666666666666664e-2);
+  cosPolynomial = fma(cosPolynomial, x2, -5.0e-1);
+  let cosine = fma(cosPolynomial, x2, 1.0);
+  return vec2<f32>(cosine, sine);
+}
+{% if needPreciseTrigCentered %}
+// Input must be in [-pi, pi]. Returns (cos(x), sin(x)).
+fn precise_sincos_centered(x: f32) -> vec2<f32> {
+  var folded = x;
+  var cosineSign = 1.0;
+  if (folded > PRECISE_TRIG_HALF_PI) {
+    folded = PRECISE_TRIG_PI - folded;
+    cosineSign = -1.0;
+  } else if (folded < -PRECISE_TRIG_HALF_PI) {
+    folded = -PRECISE_TRIG_PI - folded;
+    cosineSign = -1.0;
+  }
+  let value = precise_sincos_half_pi(folded);
+  return vec2<f32>(cosineSign * value.x, value.y);
+}
+{% endif %}
+// 24-bit words of 2/pi used by the exact Payne-Hanek large-input reducer.
+const TWO_OVER_PI: array<u32, 14> = array<u32, 14>(
+  0xa2f983u, 0x6e4e44u, 0x1529fcu, 0x2757d1u, 0xf534ddu, 0xc0db62u,
+  0x95993cu, 0x439041u, 0xfe5163u, 0xabdebbu, 0xc561b7u, 0x246e3au,
+  0x424dd2u, 0xe00649u
+);
+const PIO2_F: f32 = 1.5707963267948966;
+const INV_PIO2_F: f32 = 0.6366197723675814;
+const PIO2_HI_F: f32 = 1.570796251296997;
+const PIO2_LO_F: f32 = 7.549789415861596e-8;
+const REDUCE_THRESHOLD: f32 = 1.0e4;
+struct Pio2 { octant: u32, r: f32 };
+{{ emit_reduce_pio2_fast() }}
+{{ emit_reduce_pio2() }}
+{%- endmacro %}
+{% macro emit_cos_accurate() %}
+fn cos_accurate(x: f32) -> f32 {
+  let ax = abs(x);                 // cos is even
+  let ax_bits = bitcast<u32>(ax);
+  if (((ax_bits >> 23u) & 0xffu) == 0xffu) {
+    // Avoid Inf-Inf / NaN-NaN arithmetic: setting the quiet-NaN bit is robust
+    // even on backends that compile arithmetic under a no-NaN assumption.
+    return bitcast<f32>(ax_bits | 0x00400000u);
+  }
+  var red: Pio2;
+  if (ax < REDUCE_THRESHOLD) { red = reduce_pio2_fast(ax); }
+  else { red = reduce_pio2(ax); }
+  let reduced = precise_sincos_half_pi(red.r);
+  switch (red.octant) {
+    case 0u: { return reduced.x; }
+    case 1u: { return -reduced.y; }
+    case 2u: { return -reduced.x; }
+    default: { return reduced.y; }
+  }
+}{% endmacro %}
+{{ emit_trig_reduction_support() }}
+{{ emit_cos_accurate() }}
+{{ flat_tail_open() }}
+  y[i] = {{ scalar }}(cos_accurate(f32(x[i])));
+{{ flat_tail_close() -}}
+}

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

	@@ -0,0 +1,198 @@

+// 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_reduce_pio2() %}
+fn reduce_pio2(ax: f32) -> Pio2 {
+  let ix = bitcast<u32>(ax);
+  let e = (ix >> 23u) & 0xffu;
+  let mant = (ix & 0x7fffffu) | 0x800000u;   // 24-bit mantissa
+  let p = i32(e) - 150;                        // ax = mant * 2^p
+  // 96-bit accumulator a2:a1:a0 holding (ax * 2/pi  mod 4) * 2^94.
+  var a0 = 0u; var a1 = 0u; var a2 = 0u;
+  for (var k = 0u; k < 14u; k = k + 1u) {
+    let shift = p - 24 * (i32(k) + 1) + 94;
+    if (shift <= -48 || shift >= 96) { continue; }
+    // prod = mant * TWO_OVER_PI[k]  (<= 48 bits) -> phi32:plo
+    let t = TWO_OVER_PI[k];
+    let mlo = mant & 0xffffu; let mhi = mant >> 16u;
+    let tlo = t & 0xffffu;    let thi = t >> 16u;
+    let ll = mlo * tlo;
+    let mid = mlo * thi + mhi * tlo;            // <= 25 bits, no overflow
+    let hh = mhi * thi;
+    let plo = ll + ((mid & 0xffffu) << 16u);
+    let carry = select(0u, 1u, plo < ll);
+    let phi32 = hh + (mid >> 16u) + carry;      // <= 16 bits
+    // Place (phi32:plo) << shift as a 64-bit value (v1:v0) at offset off >= 0.
+    var v0 = 0u; var v1 = 0u; var off = 0;
+    if (shift < 0) {
+      let s = u32(-shift);                      // 1..47
+      if (s < 32u) {
+        v0 = (plo >> s) | (phi32 << (32u - s));
+        v1 = phi32 >> s;
+      } else {
+        v0 = phi32 >> (s - 32u);
+        v1 = 0u;
+      }
+      off = 0;
+    } else {
+      v0 = plo; v1 = phi32; off = shift;        // 0..95
+    }
+    // Spread (v1:v0) across 96-bit limbs starting at bit `off`.
+    let li = off / 32; let b = u32(off % 32);
+    var w0 = 0u; var w1 = 0u; var w2 = 0u;
+    if (b == 0u) { w0 = v0; w1 = v1; w2 = 0u; }
+    else {
+      w0 = v0 << b;
+      w1 = (v1 << b) | (v0 >> (32u - b));
+      w2 = v1 >> (32u - b);
+    }
+    var t0 = 0u; var t1 = 0u; var t2 = 0u;
+    if (li == 0) { t0 = w0; t1 = w1; t2 = w2; }
+    else if (li == 1) { t1 = w0; t2 = w1; }
+    else { t2 = w0; }
+    // 96-bit add acc += t (bits above 2^96 are multiples of 4, so discarded).
+    let n0 = a0 + t0; let c0 = select(0u, 1u, n0 < a0);
+    let n1a = a1 + t1; let c1a = select(0u, 1u, n1a < a1);
+    let n1 = n1a + c0; let c1b = select(0u, 1u, n1 < n1a);
+    let n2 = a2 + t2 + c1a + c1b;
+    a0 = n0; a1 = n1; a2 = n2;
+  }
+  var n = a2 >> 30u;                            // integer part mod 4
+  var phi = f32(a2 & 0x3fffffffu) * (1.0 / 1073741824.0)
+          + f32(a1) * (1.0 / 4611686018427387904.0);   // fraction in [0,1)
+  if (phi >= 0.5) { phi = phi - 1.0; n = n + 1u; }
+  var out: Pio2;
+  out.octant = n & 3u;
+  out.r = phi * PIO2_F;
+  return out;
+}{% endmacro %}
+{% macro emit_reduce_pio2_fast() %}
+fn reduce_pio2_fast(ax: f32) -> Pio2 {
+  // Cody-Waite reduction for ordinary magnitudes. Splitting pi/2 keeps the
+  // residual accurate even when n is thousands; larger inputs use Payne-Hanek.
+  let nFloat = floor(fma(ax, INV_PIO2_F, 0.5));
+  var residual = fma(-nFloat, PIO2_HI_F, ax);
+  residual = fma(-nFloat, PIO2_LO_F, residual);
+  var out: Pio2;
+  out.octant = u32(nFloat) & 3u;
+  out.r = residual;
+  return out;
+}{% endmacro %}
+{% macro emit_trig_reduction_support() %}
+{% set needPreciseTrigCentered = false %}
+{% set needPreciseTrigTwoPi = false %}
+// Backend-stable f32 sine/cosine core.
+//
+// Shader transcendental accuracy is implementation-defined, and some portable
+// backends are only accurate to roughly 1e-4. Each path retains the most accurate
+// available phase representation, reduces it to [-pi, pi], then uses these
+// polynomials. The half-pi core is degree 13 for sine and degree 12 for cosine;
+// truncation error is well below one f32 ULP over its documented interval.
+{% if needPreciseTrigCentered %}
+const PRECISE_TRIG_PI: f32 = 3.141592653589793;
+{% endif %}
+{% if needPreciseTrigTwoPi %}
+const PRECISE_TRIG_TWO_PI: f32 = 6.283185307179586;
+{% endif %}
+{% if needPreciseTrigCentered %}
+const PRECISE_TRIG_HALF_PI: f32 = 1.5707963267948966;
+{% endif %}
+// Input must be in [-pi/2, pi/2]. Returns (cos(x), sin(x)).
+fn precise_sincos_half_pi(x: f32) -> vec2<f32> {
+  let x2 = x * x;
+  var sinPolynomial = 1.6059043836821613e-10;
+  sinPolynomial = fma(sinPolynomial, x2, -2.505210838544172e-8);
+  sinPolynomial = fma(sinPolynomial, x2, 2.7557319223985893e-6);
+  sinPolynomial = fma(sinPolynomial, x2, -1.984126984126984e-4);
+  sinPolynomial = fma(sinPolynomial, x2, 8.333333333333333e-3);
+  sinPolynomial = fma(sinPolynomial, x2, -1.6666666666666666e-1);
+  let sine = x * fma(sinPolynomial, x2, 1.0);
+  var cosPolynomial = 2.08767569878681e-9;
+  cosPolynomial = fma(cosPolynomial, x2, -2.755731922398589e-7);
+  cosPolynomial = fma(cosPolynomial, x2, 2.48015873015873e-5);
+  cosPolynomial = fma(cosPolynomial, x2, -1.388888888888889e-3);
+  cosPolynomial = fma(cosPolynomial, x2, 4.1666666666666664e-2);
+  cosPolynomial = fma(cosPolynomial, x2, -5.0e-1);
+  let cosine = fma(cosPolynomial, x2, 1.0);
+  return vec2<f32>(cosine, sine);
+}
+{% if needPreciseTrigCentered %}
+// Input must be in [-pi, pi]. Returns (cos(x), sin(x)).
+fn precise_sincos_centered(x: f32) -> vec2<f32> {
+  var folded = x;
+  var cosineSign = 1.0;
+  if (folded > PRECISE_TRIG_HALF_PI) {
+    folded = PRECISE_TRIG_PI - folded;
+    cosineSign = -1.0;
+  } else if (folded < -PRECISE_TRIG_HALF_PI) {
+    folded = -PRECISE_TRIG_PI - folded;
+    cosineSign = -1.0;
+  }
+  let value = precise_sincos_half_pi(folded);
+  return vec2<f32>(cosineSign * value.x, value.y);
+}
+{% endif %}
+// 24-bit words of 2/pi used by the exact Payne-Hanek large-input reducer.
+const TWO_OVER_PI: array<u32, 14> = array<u32, 14>(
+  0xa2f983u, 0x6e4e44u, 0x1529fcu, 0x2757d1u, 0xf534ddu, 0xc0db62u,
+  0x95993cu, 0x439041u, 0xfe5163u, 0xabdebbu, 0xc561b7u, 0x246e3au,
+  0x424dd2u, 0xe00649u
+);
+const PIO2_F: f32 = 1.5707963267948966;
+const INV_PIO2_F: f32 = 0.6366197723675814;
+const PIO2_HI_F: f32 = 1.570796251296997;
+const PIO2_LO_F: f32 = 7.549789415861596e-8;
+const REDUCE_THRESHOLD: f32 = 1.0e4;
+struct Pio2 { octant: u32, r: f32 };
+{{ emit_reduce_pio2_fast() }}
+{{ emit_reduce_pio2() }}
+{%- endmacro %}
+{% macro emit_cos_accurate() %}
+fn cos_accurate(x: f32) -> f32 {
+  let ax = abs(x);                 // cos is even
+  let ax_bits = bitcast<u32>(ax);
+  if (((ax_bits >> 23u) & 0xffu) == 0xffu) {
+    // Avoid Inf-Inf / NaN-NaN arithmetic: setting the quiet-NaN bit is robust
+    // even on backends that compile arithmetic under a no-NaN assumption.
+    return bitcast<f32>(ax_bits | 0x00400000u);
+  }
+  var red: Pio2;
+  if (ax < REDUCE_THRESHOLD) { red = reduce_pio2_fast(ax); }
+  else { red = reduce_pio2(ax); }
+  let reduced = precise_sincos_half_pi(red.r);
+  switch (red.octant) {
+    case 0u: { return reduced.x; }
+    case 1u: { return -reduced.y; }
+    case 2u: { return -reduced.x; }
+    default: { return reduced.y; }
+  }
+}{% endmacro %}
+{{ emit_trig_reduction_support() }}
+{{ emit_cos_accurate() }}
+@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>(cos_accurate(fv.x), cos_accurate(fv.y), cos_accurate(fv.z), cos_accurate(fv.w)));
+}